[FEA]: Support Unary Map on DeviceReduce and use that from Thrust reduce
- Dominant language
- C++
- Stars
- 2.5k
- Forks
- 487
- Avg merge
- 2d 7h
- Merged PRs (30d)
- 296
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
General CCCL
### Is your feature request related to a problem? Please describe.
Currently cub DeviceReduce only supports a BinaryReduceOp , and to perform a Map-Reduce, one needs to use a TransformIterator, but unfortunately doing so prevent load vectorization within DeviceReduce.
In my experiment the impact of this on the performance of thrust::transform_reduce is about 6x in some cases.
### Describe the solution you'd like
The solution I've implemented is to support a UnaryMapOp in DeviceReduce that defaults to an Identity mapping function, avoiding having to duplicate code and interfaces, but which may be an API breaking change.
### Describe alternatives you've considered
_No response_
### Additional context
_No response_
Contributor guide
Research direction
Start at the DeviceReduce entry point and the thrust::transform_reduce path described in the issue, then investigate how TransformIterator currently affects load vectorization. Done means DeviceReduce accepts an optional UnaryMapOp with identity behavior and Thrust reduce uses it without the reported vectorization regression; check API compatibility and performance against the stated experiment.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- cpp
- Domain
- hpc
- Issue type
- Feature
- Difficulty
- 5/5
- Estimated time
- Over a week
- Activity status
- Stale
- Clarity
- Mostly clear
- Newbie friendliness
- 35/100