[FEA]: Introduce API-Entry-Level Error Checks
- Dominant language
- C++
- Stars
- 2.5k
- Forks
- 486
- Avg merge
- 2d 6h
- Merged PRs (30d)
- 295
Description
### Is this a duplicate?
- [x] I confirmed there appear to be no [duplicate issues](https://github.com/NVIDIA/cccl/issues) for this request and that I agree to the [Code of Conduct](CODE_OF_CONDUCT.md)
### Area
CUB
### Is your feature request related to a problem? Please describe.
@elstehle illustrated a potential error masking issue:
```cpp
produce_async_oob_access();
cudaError_t error = cub::Scan(d_temp_storage, wrong_temp_storage_bytes); // cudaErrorInvalidValue
```
In the example above, `CubDebug` around `AliasTemporaries` would reset asynchronous error (if observed) and return `cudaErrorInvalidValue`. Any CUB error would be preferred to the one returned by `cudaGetLastError` inside `CubDebug`.
### Describe the solution you'd like
With a sufficient `CUB_DEBUG_*` level we synchronize our kernels to catch potential issues. We could make user lives easier if we did the same at the entry of each device-scope algorithm:
```cpp
error = detail::DebugSyncStream(stream);
if (CubDebug(error))
{
break;
}
```
### Describe alternatives you've considered
We considered resetting errors only when `CubDebug` is given `cudaSuccess`. The downside of this solution is that we'd enforce users to reset errors after each CUB invocation. I incline towards CUB algorithms having `cudaGetLastError` semantics, if the error is observed, it's reset.
### Additional context
_No response_
Contributor guide
Assessment
This issue has not been assessed yet.