[FEA]: Redesign default tuning
- 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 of the parameters in parallel algorithms affect performance characteristics but not functional correctness.
Thread block size, grain size (number of items per thread), cache modifiers etc. are all examples of such parameters.
A set of these parameters composes a **tuning**:
https://github.com/NVIDIA/cccl/blob/28974d04b8b18a167e721b0abd8e24b2ac8cdd15/cub/cub/device/dispatch/tuning/tuning_scan.cuh#L216-L219
Today, we are tuning parallel algorithms for specific compile-time workloads.
For example, prefix sum of `int64_t` with `cuda::std::plus` will be tuned differently compared to prefix sum of `uint8_t` with `cuda::std::minimum`.
This approach leads to CUB not applying any of the tunings when, say, binary operator or the value types are not known.
This leads to suboptimal performance in cases that represent a slight deviation from trivial binary operators that we know about.
Initially, this approach was chosen as a safe default.
We can't introspect the incoming binary operator.
If this binary operator is complex enough, or leads to significant load imbalance, increased block and grain sizes that suit trivial cases can lead to performance regressions.
Our current intuition is that arithmetically intense operators are rare compared to trivial ones. For instance:
1. https://github.com/NVIDIA/cccl/blob/0b5844faf7b3d7e7dc47b66f756a3ef2ebfb48ac/cub/benchmarks/bench/transform/babelstream2.cu#L44-L46
2. https://github.com/NVIDIA/cccl/blob/0b5844faf7b3d7e7dc47b66f756a3ef2ebfb48ac/cub/benchmarks/bench/transform/babelstream3.cu#L36-L38
In the cases above, we'd prefer to have some tuning instead of pessimizing performance.
This issue is meant as a discussion / design point on our options in avoiding this pessimization.
### Describe the solution you'd like
We should revisit our approach on tuning unknown workloads.
There are a few options that we have considered.
## Opt-In Scheme
We could give users an opt-in mechanism to let CUB know that given value type / operator combination are trivial:
```cuda
cub::DeviceReduce(...,
proclaim_trivial_case(
[]__host__ __device__(char a, char b) -> char { return a + b }
));
```
Then, on the tuning end, we'd relax tuning from:
```c++
template
struct sm90_tuning : sm90_tuning_vals {};
```
to something like:
```c++
template
struct sm90_tuning : sm90_tuning_vals {};
```
### Pros
- User gains access to the tunings we have for trivial cases
- Heavy-weight operations do not experience performance regressions
- Annotation can propagate through the stack. Thrust user can annotate a function and this annotation would go all the way to CUB.
### Cons
- Manual annotation means that some users will leave this option on the table
## Opt-Out Scheme
Alternatively, we could apply existing tunings by default and give users a way to opt-out of these tunings:
```c++
cub::DeviceReduce(...,
proclaim_heavy_weight_case(
[]__host__ __device__(cuda::std::complex a, cuda::std::complex b) -> char {
return heavy_weight_comparison(a, b) ? a : b; }
));
```
### Pros
- Common use cases will likely experience performance improvements
- We'll be able to recognize heavy-weight workloads, and, say, remove `#pragma unroll` from CUB kernels
- We'll be able to go beyond tuning, and select completely different (work-efficient) implementations of parallel algorithms for this case
- Annotation can propagate through the stack. Thrust user can annotate a function and this annotation would go all the way to CUB.
### Cons
- Users providing heavy-weight operations will likely experience perf regressions and will have to manually annotate their code
I incline towards opt-out scheme.
Before closing this issue, we should allow for some time to gather more opinions.
After that, the issue can be closed by a design of function wrapper with an example of opt-in / opt-out (depending on what we choose) behavior in one of CUB algorithms.
### Describe alternatives you've considered
_No response_
### Additional context
_No response_
Contributor guide
Assessment
This issue has not been assessed yet.