NVIDIA / NVIDIA/cccl

[FEA]: Add `DevicePartition::StencilIf` interface and add `DevicePartition::{If,Flagged,StencilIf}` overloads that take two distinct iterators

Open
#1,436 0 comments 0 reactions 0 assignees 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.

**`DevicePartition::StencilIf`**
As part of https://github.com/NVIDIA/cccl/pull/1379, the agent (`AgentSelectIf`) that implements both the `DeviceSelect::{If,Flagged}` algorithms as well as `DevicePartition::{If,Flagged}` algorithms, was extended to support `StencilIf` operations. This capability of taking both a `flagged` (aka `stencil`) iterator _as well as_ the `select` callable is theoretically available to both `DeviceSelect` and `DevicePartition`, but not yet exposed in the `DevicePartition` interfaces.

Hence, similar to https://github.com/NVIDIA/cccl/issues/1409, we should add a `DevicePartition::StencilIf` interface

**Adding overloads to partition to two distinct iterators**
As part of https://github.com/NVIDIA/cccl/pull/1435, we added the option for `cub::DevicePartition` to partition the inputs to two distinct iterators. That is, for two user-provided iterators `selected_it` and `rejected_it`, all selected items are written to `selected_it` and all rejected items are written to `rejected_it`.

This can be achieved by wrapping the two output iterators in the `cub::detail::partition_distinct_output_t{selected_it, rejected_it}` wrapper struct and passing that as the output iterator to `DispatchSelectIf`. The agent is specialized to scatter the items to the two iterators wrapped by the `partition_distinct_output_t` instance.

We want _add_ overloads to `DevicePartition::{If,Flagged,StencilIf}` that take two distinct iterators for selected and rejected items, respectively.

### Tasks
- [ ] Add new interface for `cub::DevicePartition::StencilIf`
- [ ] Add documentation and example for `cub::DeviceSelect::StencilIf`
- [ ] Add tests for `cub::DeviceSelect::StencilIf`
- [ ] Add overloads to `DevicePartition::{If,Flagged,StencilIf}` that take two distinct iterators for selected and rejected items, respectively.
- [ ] Add documentation and examples for partitioning to distinct iterators
- [ ] Add tests for the new overload

### Describe the solution you'd like

Expose new functionality via the new `DevicePartition` interfaces.

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