[BUG]: Many cuco functions incorrectly marked as `noexcept`

[wence-]

Is this a duplicate?

• I confirmed there appear to be no duplicate issues for this bug (https://github.com/NVIDIA/cuCollections/issues)

Type of Bug

Runtime Error

Describe the bug

Many cuco functions, both ctors and methods, are incorrectly marked as noexcept, despite them making calls into the cuda runtime that might potentially fail.

Some (non-exhaustive) examples:

• aow_storage::initialize, calls cuda_stream_ref::synchronize which might throw
• aow_storage ctor, allocates memory which might throw
• open_addressing_impl ctor, allocates memory which might throw
• open_addressing_impl::size calls a kernel and syncs the stream which might throw
• open_addressing_impl::initialize calls aow_storage::initialize and therefore transitively might throw
• distinct_count_estimator::add_async is marked as noexcept, but the impl class's method (hyperloglog::add_async) is not noexcept, and indeed might throw since it launches kernels and calls CUCO_CUDA_TRY(cudaLaunchKernel(...)).
• static_map::clear is marked noexcept, but open_addressing_impl::clear calls aow_storage::initialize and therefore transitively might throw

How to Reproduce

Try and construct a static set requiring more storage than there is device memory (or perhaps using RMM with a limited memory resource). RMM will throw a subclass of std::bad_alloc but we will not be able to catch it with:

try {
   make_cuco_static_set(...);
} catch (...) {
   // expect to get here
}

and instead the program will abort.

Expected behavior

It should be possible to catch exceptions thrown inside cuco functions.

Reproduction link

No response

Operating System

No response

nvidia-smi output

No response

NVCC version

No response

[wence-]

See rapidsai/cudf#16059 for an example of this.

[sleeepyjack]

Good catch! I'll take a look and open a PR to fix this. Maybe we should also add some diagnostics to CI to catch such errors automatically but it seems that it hasn't been implemented yet at least in gcc: https://gcc.gnu.org/bugzilla/show_bug.cgi?id=61372

Any idea for a workaround?

[wence-]

Good catch! I'll take a look and open a PR to fix this. Maybe we should also add some diagnostics to CI to catch such errors automatically but it seems that it hasn't been implemented yet at least in gcc: https://gcc.gnu.org/bugzilla/show_bug.cgi?id=61372

Any idea for a workaround?

clang-tidy offers bugprone-exception-escape but I'm not sure if one can use clang-tidy with cuda code.

[PointKernel]

I'm not sure if one can use clang-tidy with cuda code.

As long as the device code can be built with clang, clang-tidy will work fine.

#129 kicked off this effort but was blocked by libcudacxx and nvbench couldn't be built with clang at that point.