Is this a duplicate?

• I confirmed there appear to be no duplicate issues for this request and that I agree to the Code of Conduct

Area

CUB

Is your feature request related to a problem? Please describe.

The cub::DeviceReduce::Arg{Min,Max} interfaces currently return the (extremum, index)-tuple of the result as a cub::KeyValuePair. This interface has shortcomings that were discussed in #414 (comment) and in https://github.com/NVIDIA/cccl/pull/2647/files#r1824185104.

Describe the solution you'd like

The following interface works around the referenced shortcomings and provides users with a more flexible and convenient interface, allowing to mix-and-match arbitrary output iterator combinations to which the index and extremum are written:

static cudaError_t ArgMin(
    void* d_temp_storage,
    size_t& temp_storage_bytes,
    InputIteratorT d_in,
    ExtremumOutIteratorT d_min_out,
    IndexOutIteratorT d_index_out,
    ::cuda::std::int64_t num_items,
    cudaStream_t stream)

Describe alternatives you've considered

No response

Additional context

No response