Skip to content

Commit d36905c

Browse files
PointKernelPointKernel
authored
Get rid of the stream_ref default construction to avoid deprecation warnings (#763)
The default constructor of `cuda::stream_ref` is being deprecated in CCCL, and this PR addresses the warnings by replacing default construction of `cuda::stream_ref`, e.g. https://godbolt.org/z/rd45Mx9dq: ```bash warning: 'constexpr cuda::__4::stream_ref::stream_ref()' is deprecated: Using the default/null stream is generally discouraged. If you need to use it, please construct a stream_ref from cudaStream_t{nullptr} ```
1 parent 2777267 commit d36905c

16 files changed

+382
-214
lines changed

‎include/cuco/bloom_filter.cuh‎

Lines changed: 26 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
/*
2-
* Copyright (c) 2024, NVIDIA CORPORATION.
2+
* Copyright (c) 2024-2025, NVIDIA CORPORATION.
33
*
44
* Licensed under the Apache License, Version 2.0 (the "License");
55
* you may not use this file except in compliance with the License.
@@ -117,7 +117,8 @@ class bloom_filter {
117117
cuda_thread_scope<Scope> scope = {},
118118
Policy const& policy = {},
119119
Allocator const& alloc = {},
120-
cuda::stream_ref stream = {});
120+
cuda::stream_ref stream = cuda::stream_ref{
121+
cudaStream_t{nullptr}});
121122

122123
/**
123124
* @brief Erases all information from the filter.
@@ -127,14 +128,15 @@ class bloom_filter {
127128
*
128129
* @param stream CUDA stream used for device memory operations and kernel launches
129130
*/
130-
__host__ constexpr void clear(cuda::stream_ref stream = {});
131+
__host__ constexpr void clear(cuda::stream_ref stream = cuda::stream_ref{cudaStream_t{nullptr}});
131132

132133
/**
133134
* @brief Asynchronously erases all information from the filter.
134135
*
135136
* @param stream CUDA stream used for device memory operations and kernel launches
136137
*/
137-
__host__ constexpr void clear_async(cuda::stream_ref stream = {});
138+
__host__ constexpr void clear_async(cuda::stream_ref stream = cuda::stream_ref{
139+
cudaStream_t{nullptr}});
138140

139141
/**
140142
* @brief Adds all keys in the range `[first, last)` to the filter.
@@ -148,7 +150,9 @@ class bloom_filter {
148150
* @param stream CUDA stream used for device memory operations and kernel launches
149151
*/
150152
template <class InputIt>
151-
__host__ constexpr void add(InputIt first, InputIt last, cuda::stream_ref stream = {});
153+
__host__ constexpr void add(InputIt first,
154+
InputIt last,
155+
cuda::stream_ref stream = cuda::stream_ref{cudaStream_t{nullptr}});
152156

153157
/**
154158
* @brief Asynchronously adds all keys in the range `[first, last)` to the filter.
@@ -160,7 +164,8 @@ class bloom_filter {
160164
* @param stream CUDA stream used for device memory operations and kernel launches
161165
*/
162166
template <class InputIt>
163-
__host__ constexpr void add_async(InputIt first, InputIt last, cuda::stream_ref stream = {});
167+
__host__ constexpr void add_async(
168+
InputIt first, InputIt last, cuda::stream_ref stream = cuda::stream_ref{cudaStream_t{nullptr}});
164169

165170
/**
166171
* @brief Adds keys in the range `[first, last)` if `pred` of the corresponding `stencil` returns
@@ -184,8 +189,11 @@ class bloom_filter {
184189
* @param stream CUDA stream used for device memory operations and kernel launches
185190
*/
186191
template <class InputIt, class StencilIt, class Predicate>
187-
__host__ constexpr void add_if(
188-
InputIt first, InputIt last, StencilIt stencil, Predicate pred, cuda::stream_ref stream = {});
192+
__host__ constexpr void add_if(InputIt first,
193+
InputIt last,
194+
StencilIt stencil,
195+
Predicate pred,
196+
cuda::stream_ref stream = cuda::stream_ref{cudaStream_t{nullptr}});
189197

190198
/**
191199
* @brief Asynchronously adds keys in the range `[first, last)` if `pred` of the corresponding
@@ -211,7 +219,8 @@ class bloom_filter {
211219
InputIt last,
212220
StencilIt stencil,
213221
Predicate pred,
214-
cuda::stream_ref stream = {}) noexcept;
222+
cuda::stream_ref stream = cuda::stream_ref{
223+
cudaStream_t{nullptr}}) noexcept;
215224

216225
/**
217226
* @brief Tests all keys in the range `[first, last)` if their fingerprints are present in the
@@ -232,7 +241,8 @@ class bloom_filter {
232241
__host__ constexpr void contains(InputIt first,
233242
InputIt last,
234243
OutputIt output_begin,
235-
cuda::stream_ref stream = {}) const;
244+
cuda::stream_ref stream = cuda::stream_ref{
245+
cudaStream_t{nullptr}}) const;
236246

237247
/**
238248
* @brief Asynchronously tests all keys in the range `[first, last)` if their fingerprints are
@@ -250,7 +260,8 @@ class bloom_filter {
250260
__host__ constexpr void contains_async(InputIt first,
251261
InputIt last,
252262
OutputIt output_begin,
253-
cuda::stream_ref stream = {}) const noexcept;
263+
cuda::stream_ref stream = cuda::stream_ref{
264+
cudaStream_t{nullptr}}) const noexcept;
254265

255266
/**
256267
* @brief Tests all keys in the range `[first, last)` if their fingerprints are present in the
@@ -281,7 +292,8 @@ class bloom_filter {
281292
StencilIt stencil,
282293
Predicate pred,
283294
OutputIt output_begin,
284-
cuda::stream_ref stream = {}) const;
295+
cuda::stream_ref stream = cuda::stream_ref{
296+
cudaStream_t{nullptr}}) const;
285297

286298
/**
287299
* @brief Asynchronously tests all keys in the range `[first, last)` if their fingerprints are
@@ -310,7 +322,8 @@ class bloom_filter {
310322
StencilIt stencil,
311323
Predicate pred,
312324
OutputIt output_begin,
313-
cuda::stream_ref stream = {}) const noexcept;
325+
cuda::stream_ref stream = cuda::stream_ref{
326+
cudaStream_t{nullptr}}) const noexcept;
314327

315328
/**
316329
* @brief Gets a pointer to the underlying filter storage.

‎include/cuco/bloom_filter_ref.cuh‎

Lines changed: 23 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -101,14 +101,15 @@ class bloom_filter_ref {
101101
*
102102
* @param stream CUDA stream used for device memory operations and kernel launches
103103
*/
104-
__host__ constexpr void clear(cuda::stream_ref stream = {});
104+
__host__ constexpr void clear(cuda::stream_ref stream = cuda::stream_ref{cudaStream_t{nullptr}});
105105

106106
/**
107107
* @brief Asynchronously erases all information from the filter.
108108
*
109109
* @param stream CUDA stream used for device memory operations and kernel launches
110110
*/
111-
__host__ constexpr void clear_async(cuda::stream_ref stream = {});
111+
__host__ constexpr void clear_async(cuda::stream_ref stream = cuda::stream_ref{
112+
cudaStream_t{nullptr}});
112113

113114
/**
114115
* @brief Device function that adds a key to the filter.
@@ -163,7 +164,9 @@ class bloom_filter_ref {
163164
* @param stream CUDA stream used for device memory operations and kernel launches
164165
*/
165166
template <class InputIt>
166-
__host__ constexpr void add(InputIt first, InputIt last, cuda::stream_ref stream = {});
167+
__host__ constexpr void add(InputIt first,
168+
InputIt last,
169+
cuda::stream_ref stream = cuda::stream_ref{cudaStream_t{nullptr}});
167170

168171
/**
169172
* @brief Asynchronously adds all keys in the range `[first, last)` to the filter.
@@ -175,7 +178,8 @@ class bloom_filter_ref {
175178
* @param stream CUDA stream used for device memory operations and kernel launches
176179
*/
177180
template <class InputIt>
178-
__host__ constexpr void add_async(InputIt first, InputIt last, cuda::stream_ref stream = {});
181+
__host__ constexpr void add_async(
182+
InputIt first, InputIt last, cuda::stream_ref stream = cuda::stream_ref{cudaStream_t{nullptr}});
179183

180184
/**
181185
* @brief Adds keys in the range `[first, last)` if `pred` of the corresponding `stencil` returns
@@ -199,8 +203,11 @@ class bloom_filter_ref {
199203
* @param stream CUDA stream used for device memory operations and kernel launches
200204
*/
201205
template <class InputIt, class StencilIt, class Predicate>
202-
__host__ constexpr void add_if(
203-
InputIt first, InputIt last, StencilIt stencil, Predicate pred, cuda::stream_ref stream = {});
206+
__host__ constexpr void add_if(InputIt first,
207+
InputIt last,
208+
StencilIt stencil,
209+
Predicate pred,
210+
cuda::stream_ref stream = cuda::stream_ref{cudaStream_t{nullptr}});
204211

205212
/**
206213
* @brief Asynchronously adds keys in the range `[first, last)` if `pred` of the corresponding
@@ -226,7 +233,8 @@ class bloom_filter_ref {
226233
InputIt last,
227234
StencilIt stencil,
228235
Predicate pred,
229-
cuda::stream_ref stream = {}) noexcept;
236+
cuda::stream_ref stream = cuda::stream_ref{
237+
cudaStream_t{nullptr}}) noexcept;
230238

231239
/**
232240
* @brief Device function that tests if a key's fingerprint is present in the filter.
@@ -283,7 +291,8 @@ class bloom_filter_ref {
283291
__host__ constexpr void contains(InputIt first,
284292
InputIt last,
285293
OutputIt output_begin,
286-
cuda::stream_ref stream = {}) const;
294+
cuda::stream_ref stream = cuda::stream_ref{
295+
cudaStream_t{nullptr}}) const;
287296

288297
/**
289298
* @brief Asynchronously tests all keys in the range `[first, last)` if their fingerprints are
@@ -303,7 +312,8 @@ class bloom_filter_ref {
303312
__host__ constexpr void contains_async(InputIt first,
304313
InputIt last,
305314
OutputIt output_begin,
306-
cuda::stream_ref stream = {}) const noexcept;
315+
cuda::stream_ref stream = cuda::stream_ref{
316+
cudaStream_t{nullptr}}) const noexcept;
307317

308318
/**
309319
* @brief Tests all keys in the range `[first, last)` if their fingerprints are present in the
@@ -336,7 +346,8 @@ class bloom_filter_ref {
336346
StencilIt stencil,
337347
Predicate pred,
338348
OutputIt output_begin,
339-
cuda::stream_ref stream = {}) const;
349+
cuda::stream_ref stream = cuda::stream_ref{
350+
cudaStream_t{nullptr}}) const;
340351

341352
/**
342353
* @brief Asynchronously tests all keys in the range `[first, last)` if their fingerprints are
@@ -367,7 +378,8 @@ class bloom_filter_ref {
367378
StencilIt stencil,
368379
Predicate pred,
369380
OutputIt output_begin,
370-
cuda::stream_ref stream = {}) const noexcept;
381+
cuda::stream_ref stream = cuda::stream_ref{
382+
cudaStream_t{nullptr}}) const noexcept;
371383

372384
/**
373385
* @brief Gets a pointer to the underlying filter storage.

‎include/cuco/bucket_storage.cuh‎

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -206,15 +206,17 @@ class bucket_storage {
206206
* @param key Key to which all keys in `slots` are initialized
207207
* @param stream Stream used for executing the kernel
208208
*/
209-
void initialize(value_type key, cuda::stream_ref stream = {});
209+
void initialize(value_type key,
210+
cuda::stream_ref stream = cuda::stream_ref{cudaStream_t{nullptr}});
210211

211212
/**
212213
* @brief Asynchronously initializes each slot in the bucket storage to contain `key`.
213214
*
214215
* @param key Key to which all keys in `slots` are initialized
215216
* @param stream Stream used for executing the kernel
216217
*/
217-
void initialize_async(value_type key, cuda::stream_ref stream = {});
218+
void initialize_async(value_type key,
219+
cuda::stream_ref stream = cuda::stream_ref{cudaStream_t{nullptr}});
218220

219221
/**
220222
* @brief Gets the total number of slot buckets in the current storage.

‎include/cuco/detail/bitwise_compare.cuh‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
/*
2-
* Copyright (c) 2020-2023, NVIDIA CORPORATION.
2+
* Copyright (c) 2020-2025, NVIDIA CORPORATION.
33
*
44
* Licensed under the Apache License, Version 2.0 (the "License");
55
* you may not use this file except in compliance with the License.

‎include/cuco/detail/open_addressing/open_addressing_impl.cuh‎

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -452,7 +452,10 @@ class open_addressing_impl {
452452
* provided at construction
453453
*/
454454
template <typename InputIt, typename Ref>
455-
void erase_async(InputIt first, InputIt last, Ref container_ref, cuda::stream_ref stream = {})
455+
void erase_async(InputIt first,
456+
InputIt last,
457+
Ref container_ref,
458+
cuda::stream_ref stream = cuda::stream_ref{cudaStream_t{nullptr}})
456459
{
457460
CUCO_EXPECTS(this->empty_key_sentinel() != this->erased_key_sentinel(),
458461
"The empty key sentinel and erased key sentinel cannot be the same value.",

‎include/cuco/detail/roaring_bitmap/roaring_bitmap_impl.cuh‎

Lines changed: 6 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -61,7 +61,7 @@ class roaring_bitmap_impl<cuda::std::uint32_t> {
6161
__host__ void contains(InputIt first,
6262
InputIt last,
6363
OutputIt contained,
64-
cuda::stream_ref stream = {}) const
64+
cuda::stream_ref stream = cuda::stream_ref{cudaStream_t{nullptr}}) const
6565
{
6666
this->contains_async(first, last, contained, stream);
6767
#if CCCL_MAJOR_VERSION > 3 || (CCCL_MAJOR_VERSION == 3 && CCCL_MINOR_VERSION >= 1)
@@ -75,7 +75,8 @@ class roaring_bitmap_impl<cuda::std::uint32_t> {
7575
__host__ void contains_async(InputIt first,
7676
InputIt last,
7777
OutputIt contained,
78-
cuda::stream_ref stream = {}) const noexcept
78+
cuda::stream_ref stream = cuda::stream_ref{
79+
cudaStream_t{nullptr}}) const noexcept
7980
{
8081
if (this->empty()) {
8182
cub::DeviceTransform::Transform(
@@ -304,7 +305,7 @@ class roaring_bitmap_impl<cuda::std::uint64_t> {
304305
__host__ void contains(InputIt first,
305306
InputIt last,
306307
OutputIt contained,
307-
cuda::stream_ref stream = {}) const
308+
cuda::stream_ref stream = cuda::stream_ref{cudaStream_t{nullptr}}) const
308309
{
309310
this->contains_async(first, last, contained, stream);
310311
#if CCCL_MAJOR_VERSION > 3 || (CCCL_MAJOR_VERSION == 3 && CCCL_MINOR_VERSION >= 1)
@@ -318,7 +319,8 @@ class roaring_bitmap_impl<cuda::std::uint64_t> {
318319
__host__ void contains_async(InputIt first,
319320
InputIt last,
320321
OutputIt contained,
321-
cuda::stream_ref stream = {}) const noexcept
322+
cuda::stream_ref stream = cuda::stream_ref{
323+
cudaStream_t{nullptr}}) const noexcept
322324
{
323325
if (this->empty()) {
324326
cub::DeviceTransform::Transform(

‎include/cuco/detail/trie/dynamic_bitset/dynamic_bitset.cuh‎

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,6 @@
11
/*
2-
* SPDX-FileCopyrightText: Copyright (c) 2023 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
3-
* SPDX-License-Identifier: Apache-2.0
2+
* SPDX-FileCopyrightText: Copyright (c) 2023-2025 NVIDIA CORPORATION & AFFILIATES. All rights
3+
* reserved. SPDX-License-Identifier: Apache-2.0
44
*
55
* Licensed under the Apache License, Version 2.0 (the "License");
66
* you may not use this file except in compliance with the License.
@@ -144,7 +144,7 @@ class dynamic_bitset {
144144
constexpr void test(KeyIt keys_begin,
145145
KeyIt keys_end,
146146
OutputIt outputs_begin,
147-
cuda::stream_ref stream = {}) noexcept;
147+
cuda::stream_ref stream = cuda::stream_ref{cudaStream_t{nullptr}}) noexcept;
148148

149149
/**
150150
* @brief For any element `keys_begin[i]` in the range `[keys_begin, keys_end)`, stores total
@@ -164,7 +164,7 @@ class dynamic_bitset {
164164
constexpr void rank(KeyIt keys_begin,
165165
KeyIt keys_end,
166166
OutputIt outputs_begin,
167-
cuda::stream_ref stream = {}) noexcept;
167+
cuda::stream_ref stream = cuda::stream_ref{cudaStream_t{nullptr}}) noexcept;
168168

169169
/**
170170
* @brief For any element `keys_begin[i]` in the range `[keys_begin, keys_end)`, stores the
@@ -184,7 +184,7 @@ class dynamic_bitset {
184184
constexpr void select(KeyIt keys_begin,
185185
KeyIt keys_end,
186186
OutputIt outputs_begin,
187-
cuda::stream_ref stream = {}) noexcept;
187+
cuda::stream_ref stream = cuda::stream_ref{cudaStream_t{nullptr}}) noexcept;
188188

189189
using rank_type = cuco::experimental::detail::rank; ///< Rank type
190190

@@ -353,7 +353,7 @@ class dynamic_bitset {
353353
*
354354
* @param stream Stream to execute kernels
355355
*/
356-
constexpr void build(cuda::stream_ref stream = {}) noexcept;
356+
constexpr void build(cuda::stream_ref stream = cuda::stream_ref{cudaStream_t{nullptr}}) noexcept;
357357

358358
/**
359359
* @brief Populates rank and select indexes for true or false bits
@@ -367,7 +367,7 @@ class dynamic_bitset {
367367
thrust::device_vector<rank_type, rank_allocator_type>& ranks,
368368
thrust::device_vector<size_type, size_allocator_type>& selects,
369369
bool flip_bits,
370-
cuda::stream_ref stream = {});
370+
cuda::stream_ref stream = cuda::stream_ref{cudaStream_t{nullptr}});
371371
};
372372

373373
} // namespace detail

‎include/cuco/dynamic_map.cuh‎

Lines changed: 5 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -111,7 +111,7 @@ class dynamic_map {
111111
cuda_thread_scope<Scope> scope = {},
112112
Storage storage = {},
113113
Allocator const& alloc = {},
114-
cuda::stream_ref stream = {});
114+
cuda::stream_ref stream = cuda::stream_ref{cudaStream_t{nullptr}});
115115

116116
/**
117117
* @brief Grows the capacity of the map so there is enough space for `n` key/value pairs.
@@ -136,7 +136,9 @@ class dynamic_map {
136136
* @param stream Stream used for executing the kernels
137137
*/
138138
template <typename InputIt>
139-
void insert(InputIt first, InputIt last, cuda::stream_ref stream = {});
139+
void insert(InputIt first,
140+
InputIt last,
141+
cuda::stream_ref stream = cuda::stream_ref{cudaStream_t{nullptr}});
140142

141143
/**
142144
* @brief Indicates whether the keys in the range `[first, last)` are contained in the map.
@@ -156,7 +158,7 @@ class dynamic_map {
156158
void contains(InputIt first,
157159
InputIt last,
158160
OutputIt output_begin,
159-
cuda::stream_ref stream = {}) const;
161+
cuda::stream_ref stream = cuda::stream_ref{cudaStream_t{nullptr}}) const;
160162

161163
private:
162164
size_type size_{}; ///< Number of keys in the map

0 commit comments

Comments
 (0)