[FEA]: Add support for CUDA driver API in CUB
- 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
Assessment
This issue has not been assessed yet.