NVIDIA / NVIDIA/cccl

[FEA]: Support Unary Map on DeviceReduce and use that from Thrust reduce

Open
#288 0 comments 1 reaction 0 assignees View on GitHub
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

Open the contributing 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

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.