NVIDIA / NVIDIA/cccl

[FEA]: Introduce API-Entry-Level Error Checks

Open
#441 0 comments 0 reactions 0 assignees View on GitHub
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

Open the contributing guide

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.