NVIDIA / NVIDIA/cccl

[FEA]: Add support for CUDA driver API in CUB

Open
#638 2 comments 0 reactions 1 assignee Claimed by @jrhemstad 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.

Some applications can't or won't use the CUDA runtime API nor the official driver API headers + library due to licensing issues. Those applications either write their own loader or use something like https://github.com/CudaWrangler/cuew to get the required driver API entry points. cuew and the driver API can't coexist in the same module binary space since they both define the same symbols.

With CUBs current design it is hard to impossible to use the host-side of cub kernels which are usually the ones which work on the full device.

### Describe the solution you'd like

There are 3 different problems to solve lacking the runtime API

1.) CUB kernels have to be compiled to CUBINs and loaded with cuModuleLoadData and functions have to be found with cuModuleGetFunction
2.) Parameters have to be packed for each function. In the past I used successfully used properly packed structs for the function arguments.
3.) Kernel launch sizes have to be configured without any call into the driver.

This opens up the question, what is the minimal set of functions required for those tasks and if CUB could offer a dispatch table for those with a default initializer which can be used in case of statically linking the driver API.

With such dispatch table CUB could provide an API to 'prepare' an algorithm for future use like this

```
cub::cuda::runtime runtime = cuda::cuda::runtime::default();

cub::DeviceRadixSort::prepare deviceRadixSort(runtime);
deviceRadixSort.getRequiredStorage(..);

deviceRadixSort.dispatch(...);
```

### Describe alternatives you've considered

_No response_

### 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.