cub::DeviceSelect::Unique doesn't work with thrust::discard_iterator

[zasdfgbnm]

Sample code:

#include <cub/cub.cuh>
#include <thrust/iterator/discard_iterator.h>

int main() {
    int *p = nullptr;
    thrust::discard_iterator<int> d;
    size_t s;

    cub::DeviceSelect::Unique(nullptr, s, p, p, p, 0);    // works
    cub::DeviceSelect::Unique(nullptr, s, p, p, d, 0);    // works
    // cub::DeviceSelect::Unique(nullptr, s, p, d, p, 0);    // Error
    // cub::DeviceSelect::Unique(nullptr, s, p, d, d, 0);    // Error
}

#405 is also affected. I am not fixing this issue in #405, so the fix of this should fix cub::DeviceSelect::UniqueByKey as well.

[alliepiper]

More info: https://www.godbolt.org/z/8q9qTcxrK

/opt/compiler-explorer/libs/thrustcub/1.13.0/cub/block/specializations/../../block/../thread/../thread/thread_operators.cuh(64): error: no operator "==" matches these operands
            operand types are: const thrust::detail::any_assign == const thrust::detail::any_assign
          detected during:
            instantiation of "__nv_bool cub::Equality::operator()(const T &, const T &) const [with T=thrust::detail::any_assign]" 
(100): here
            instantiation of "__nv_bool cub::InequalityWrapper<EqualityOp>::operator()(const T &, const T &) [with EqualityOp=cub::Equality, T=thrust::detail::any_assign]" 
/opt/compiler-explorer/libs/thrustcub/1.13.0/cub/block/specializations/../../block/block_discontinuity.cuh(162): here
            instantiation of "__nv_bool cub::BlockDiscontinuity<T, BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH>::ApplyOp<FlagOp, false>::FlagT(FlagOp, const T &, const T &, int) [with T=thrust::detail::any_assign, BLOCK_DIM_X=128, BLOCK_DIM_Y=1, BLOCK_DIM_Z=1, PTX_ARCH=520, FlagOp=cub::InequalityWrapper<cub::Equality>]" 
/opt/compiler-explorer/libs/thrustcub/1.13.0/cub/block/specializations/../../block/block_discontinuity.cuh(321): here
            instantiation of "void cub::BlockDiscontinuity<T, BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH>::FlagHeads(FlagT (&)[ITEMS_PER_THREAD], T (&)[ITEMS_PER_THREAD], T (&)[ITEMS_PER_THREAD], FlagOp) [with T=thrust::detail::any_assign, BLOCK_DIM_X=128, BLOCK_DIM_Y=1, BLOCK_DIM_Z=1, PTX_ARCH=520, ITEMS_PER_THREAD=10, FlagT=int, FlagOp=cub::InequalityWrapper<cub::Equality>]" 
/opt/compiler-explorer/libs/thrustcub/1.13.0/cub/block/specializations/../../block/block_discontinuity.cuh(417): here
            instantiation of "void cub::BlockDiscontinuity<T, BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH>::FlagHeads(FlagT (&)[ITEMS_PER_THREAD], T (&)[ITEMS_PER_THREAD], FlagOp) [with T=thrust::detail::any_assign, BLOCK_DIM_X=128, BLOCK_DIM_Y=1, BLOCK_DIM_Z=1, PTX_ARCH=520, ITEMS_PER_THREAD=10, FlagT=int, FlagOp=cub::InequalityWrapper<cub::Equality>]" 
/opt/compiler-explorer/libs/thrustcub/1.13.0/cub/device/dispatch/../../agent/agent_select_if.cuh(334): here
            [ 2 instantiation contexts not shown ]
            instantiation of "OffsetT cub::AgentSelectIf<AgentSelectIfPolicyT, InputIteratorT, FlagsInputIteratorT, SelectedOutputIteratorT, SelectOpT, EqualityOpT, OffsetT, KEEP_REJECTS>::ConsumeTile<IS_LAST_TILE>(int, int, OffsetT, cub::AgentSelectIf<AgentSelectIfPolicyT, InputIteratorT, FlagsInputIteratorT, SelectedOutputIteratorT, SelectOpT, EqualityOpT, OffsetT, KEEP_REJECTS>::ScanTileStateT &) [with AgentSelectIfPolicyT=cub::DispatchSelectIf<int *, cub::NullType *, thrust::discard_iterator<int>, int *, cub::NullType, cub::Equality, int, false>::PtxSelectIfPolicyT, InputIteratorT=int *, FlagsInputIteratorT=cub::NullType *, SelectedOutputIteratorT=thrust::discard_iterator<int>, SelectOpT=cub::NullType, EqualityOpT=cub::Equality, OffsetT=int, KEEP_REJECTS=false, IS_LAST_TILE=false]" 
/opt/compiler-explorer/libs/thrustcub/1.13.0/cub/device/dispatch/../../agent/agent_select_if.cuh(681): here
            instantiation of "void cub::AgentSelectIf<AgentSelectIfPolicyT, InputIteratorT, FlagsInputIteratorT, SelectedOutputIteratorT, SelectOpT, EqualityOpT, OffsetT, KEEP_REJECTS>::ConsumeRange(int, cub::AgentSelectIf<AgentSelectIfPolicyT, InputIteratorT, FlagsInputIteratorT, SelectedOutputIteratorT, SelectOpT, EqualityOpT, OffsetT, KEEP_REJECTS>::ScanTileStateT &, NumSelectedIteratorT) [with AgentSelectIfPolicyT=cub::DispatchSelectIf<int *, cub::NullType *, thrust::discard_iterator<int>, int *, cub::NullType, cub::Equality, int, false>::PtxSelectIfPolicyT, InputIteratorT=int *, FlagsInputIteratorT=cub::NullType *, SelectedOutputIteratorT=thrust::discard_iterator<int>, SelectOpT=cub::NullType, EqualityOpT=cub::Equality, OffsetT=int, KEEP_REJECTS=false, NumSelectedIteratorT=int *]" 
/opt/compiler-explorer/libs/thrustcub/1.13.0/cub/device/dispatch/dispatch_select_if.cuh(108): here
            instantiation of "void cub::DeviceSelectSweepKernel<AgentSelectIfPolicyT,InputIteratorT,FlagsInputIteratorT,SelectedOutputIteratorT,NumSelectedIteratorT,ScanTileStateT,SelectOpT,EqualityOpT,OffsetT,KEEP_REJECTS>(InputIteratorT, FlagsInputIteratorT, SelectedOutputIteratorT, NumSelectedIteratorT, ScanTileStateT, SelectOpT, EqualityOpT, OffsetT, int) [with AgentSelectIfPolicyT=cub::DispatchSelectIf<int *, cub::NullType *, thrust::discard_iterator<int>, int *, cub::NullType, cub::Equality, int, false>::PtxSelectIfPolicyT, InputIteratorT=int *, FlagsInputIteratorT=cub::NullType *, SelectedOutputIteratorT=thrust::discard_iterator<int>, NumSelectedIteratorT=int *, ScanTileStateT=cub::ScanTileState<int, true>, SelectOpT=cub::NullType, EqualityOpT=cub::Equality, OffsetT=int, KEEP_REJECTS=false]" 
/opt/compiler-explorer/libs/thrustcub/1.13.0/cub/device/dispatch/dispatch_select_if.cuh(420): here
            instantiation of "cudaError_t cub::DispatchSelectIf<InputIteratorT, FlagsInputIteratorT, SelectedOutputIteratorT, NumSelectedIteratorT, SelectOpT, EqualityOpT, OffsetT, KEEP_REJECTS>::Dispatch(void *, size_t &, InputIteratorT, FlagsInputIteratorT, SelectedOutputIteratorT, NumSelectedIteratorT, SelectOpT, EqualityOpT, OffsetT, cudaStream_t, __nv_bool) [with InputIteratorT=int *, FlagsInputIteratorT=cub::NullType *, SelectedOutputIteratorT=thrust::discard_iterator<int>, NumSelectedIteratorT=int *, SelectOpT=cub::NullType, EqualityOpT=cub::Equality, OffsetT=int, KEEP_REJECTS=false]" 
/opt/compiler-explorer/libs/thrustcub/1.13.0/cub/device/device_select.cuh(355): here
            instantiation of "cudaError_t cub::DeviceSelect::Unique(void *, size_t &, InputIteratorT, OutputIteratorT, NumSelectedIteratorT, int, cudaStream_t, __nv_bool) [with InputIteratorT=int *, OutputIteratorT=thrust::discard_iterator<int>, NumSelectedIteratorT=int *]" 
<source>(11): here

[fkallen]

Same errors occurs here NVIDIA/cccl#784

[alliepiper]

Fixed by NVIDIA/thrust#444.