[BUG]: `cub::DeviceRunLengthEncode::Encode` does not compile when unique_output is a `thrust::transform_output_iterator`
- 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 bug and that I agree to the [Code of Conduct](CODE_OF_CONDUCT.md)
### Type of Bug
Compile-time Error
### Component
CUB
### Describe the bug
I am computing an RLE with tuple inputs, but for the output I only need one element of the tuple. To reduce the allocated memory for output, I would like to use a transform_output_iterator which maps a tuple to the required element. This does not compile.
Maybe related: https://github.com/NVIDIA/cccl/issues/821
The compilation error is very verbose.
```
/opt/compiler-explorer/cuda/13.3.0/bin/../targets/x86_64-linux/include/cccl/cub/block/block_load.cuh(117): error: no suitable conversion function from "thrust::_V_300303_SM_1200::iterator_facade>, thrust::_V_300303_SM_1200::detail::make_zip_iterator_base>::value_type, thrust::_V_300303_SM_1200::detail::make_zip_iterator_base>::system, thrust::_V_300303_SM_1200::detail::make_zip_iterator_base>::traversal_category, thrust::_V_300303_SM_1200::detail::make_zip_iterator_base>::reference, thrust::_V_300303_SM_1200::detail::make_zip_iterator_base>::difference_type>::reference" (aka "thrust::_V_300303_SM_1200::detail::tuple_of_iterator_references") to "cub::_V_300303_SM_1200::detail::reduce::AgentReduceByKey, cub::_V_300303_SM_1200::detail::non_void_value_t>, cub::_V_300303_SM_1200::detail::non_void_value_t>>, cub::_V_300303_SM_1200::detail::non_void_value_t)->cuda::std::__4::tuple_element_t>::type>, int *>, cub::_V_300303_SM_1200::detail::it_value_t>>>>::Policy1000::ReduceByKeyPolicyT, thrust::_V_300303_SM_1200::zip_iterator>, thrust::_V_300303_SM_1200::transform_output_iterator, thrust::_V_300303_SM_1200::detail::it_value_t>)->cuda::std::__4::tuple_element_t<0UL, cuda::std::__4::tuple, thrust::_V_300303_SM_1200::detail::it_value_t>>>::type>, int *>, thrust::_V_300303_SM_1200::constant_iterator>, cub::_V_300303_SM_1200::detail::choose_signed_offset_t, thrust::_V_300303_SM_1200::use_default>, int *, int *, cuda::std::__4::equal_to, cuda::std::__4::plus, cub::_V_300303_SM_1200::detail::reduce::DispatchStreamingReduceByKey>, thrust::_V_300303_SM_1200::transform_output_iterator, thrust::_V_300303_SM_1200::detail::it_value_t>)->cuda::std::__4::tuple_element_t<0UL, cuda::std::__4::tuple>>::type>, int *>, thrust::_V_300303_SM_1200::constant_iterator>, cub::_V_300303_SM_1200::detail::choose_signed_offset_t, thrust::_V_300303_SM_1200::use_default>, int *, int *, cuda::std::__4::equal_to, cuda::std::__4::plus, cub::_V_300303_SM_1200::detail::choose_signed_offset_t, cuda::std::__4::__accumulator_t, cub::_V_300303_SM_1200::detail::non_void_value_t>, cub::_V_300303_SM_1200::detail::non_void_value_t>>, cub::_V_300303_SM_1200::detail::rle::encode::policy_hub, cub::_V_300303_SM_1200::detail::non_void_value_t>, cub::_V_300303_SM_1200::detail::non_void_value_t>>, cub::_V_300303_SM_1200::detail::non_void_value_t::type>, int *>, cub::_V_300303_SM_1200::detail::it_value_t>>>>>::local_offset_t, cuda::std::__4::__accumulator_t, cub::_V_300303_SM_1200::detail::non_void_value_t>, cub::_V_300303_SM_1200::detail::non_void_value_t>>, cub::_V_300303_SM_1200::detail::reduce::DispatchStreamingReduceByKey>, thrust::_V_300303_SM_1200::transform_output_iterator, thrust::_V_300303_SM_1200::detail::it_value_t>)->cuda::std::__4::tuple_element_t<0UL, cuda::std::__4::tuple>>::type>, int *>, thrust::_V_300303_SM_1200::constant_iterator>, cub::_V_300303_SM_1200::detail::choose_signed_offset_t, thrust::_V_300303_SM_1200::use_default>, int *, int *, cuda::std::__4::equal_to, cuda::std::__4::plus, cub::_V_300303_SM_1200::detail::choose_signed_offset_t, cuda::std::__4::__accumulator_t, cub::_V_300303_SM_1200::detail::non_void_value_t>, cub::_V_300303_SM_1200::detail::non_void_value_t>>, cub::_V_300303_SM_1200::detail::rle::encode::policy_hub, cub::_V_300303_SM_1200::detail::non_void_value_t>, cub::_V_300303_SM_1200::detail::non_void_value_t>>, cub::_V_300303_SM_1200::detail::non_void_value_t::type>, int *>, cub::_V_300303_SM_1200::detail::it_value_t>>>>>::streaming_context_t>::KeyOutputT" exists
dst_items[i] = block_src_it[src_pos];
^
detected during:
instantiation of "void cub::_V_300303_SM_1200::LoadDirectBlocked(int, RandomAccessIterator, T (&)[ItemsPerThread], int) [with T=int, ItemsPerThread=14, RandomAccessIterator=thrust::_V_300303_SM_1200::zip_iterator>]" at line 874
instantiation of "void cub::_V_300303_SM_1200::BlockLoad::LoadInternal::Load(RandomAccessIterator, T (&)[ItemsPerThread], int) [with T=int, BlockDimX=256, ItemsPerThread=14, Algorithm=cub::_V_300303_SM_1200::BLOCK_LOAD_DIRECT, BlockDimY=1, BlockDimZ=1, Dummy=0, RandomAccessIterator=thrust::_V_300303_SM_1200::zip_iterator>]" at line 1247
instantiation of "void cub::_V_300303_SM_1200::BlockLoad::Load(RandomAccessIterator, T (&)[ItemsPerThread], int) [with T=int, BlockDimX=256, ItemsPerThread=14, Algorithm=cub::_V_300303_SM_1200::BLOCK_LOAD_DIRECT, BlockDimY=1, BlockDimZ=1, RandomAccessIterator=thrust::_V_300303_SM_1200::zip_iterator>]" at line 535 of /opt/compiler-explorer/cuda/13.3.0/bin/../targets/x86_64-linux/include/cccl/cub/agent/agent_reduce_by_key.cuh
instantiation of "void cub::_V_300303_SM_1200::detail::reduce::AgentReduceByKey::ConsumeTile(OffsetT, int, OffsetT, cub::_V_300303_SM_1200::detail::reduce::AgentReduceByKey::ScanTileStateT &) [with AgentReduceByKeyPolicyT=cub::_V_300303_SM_1200::AgentReduceByKeyPolicy<256, 14, cub::_V_300303_SM_1200::BLOCK_LOAD_DIRECT, cub::_V_300303_SM_1200::LOAD_CA, cub::_V_300303_SM_1200::BLOCK_SCAN_WARP_SCANS, cub::_V_300303_SM_1200::detail::exponential_backon_constructor_t<956U, 70U>>, KeysInputIteratorT=thrust::_V_300303_SM_1200::zip_iterator>, UniqueOutputIteratorT=thrust::_V_300303_SM_1200::transform_output_iterator)->int>, int *>, ValuesInputIteratorT=thrust::_V_300303_SM_1200::constant_iterator, AggregatesOutputIteratorT=int *, NumRunsOutputIteratorT=int *, EqualityOpT=cuda::std::__4::equal_to, ReductionOpT=cuda::std::__4::plus, OffsetT=int32_t, AccumT=int, StreamingContextT=cub::_V_300303_SM_1200::NullType, IS_LAST_TILE=false]" at line 769 of /opt/compiler-explorer/cuda/13.3.0/bin/../targets/x86_64-linux/include/cccl/cub/agent/agent_reduce_by_key.cuh
instantiation of "void cub::_V_300303_SM_1200::detail::reduce::AgentReduceByKey::ConsumeRange(OffsetT, cub::_V_300303_SM_1200::detail::reduce::AgentReduceByKey::ScanTileStateT &, int) [with AgentReduceByKeyPolicyT=cub::_V_300303_SM_1200::AgentReduceByKeyPolicy<256, 14, cub::_V_300303_SM_1200::BLOCK_LOAD_DIRECT, cub::_V_300303_SM_1200::LOAD_CA, cub::_V_300303_SM_1200::BLOCK_SCAN_WARP_SCANS, cub::_V_300303_SM_1200::detail::exponential_backon_constructor_t<956U, 70U>>, KeysInputIteratorT=thrust::_V_300303_SM_1200::zip_iterator>, UniqueOutputIteratorT=thrust::_V_300303_SM_1200::transform_output_iterator)->int>, int *>, ValuesInputIteratorT=thrust::_V_300303_SM_1200::constant_iterator, AggregatesOutputIteratorT=int *, NumRunsOutputIteratorT=int *, EqualityOpT=cuda::std::__4::equal_to, ReductionOpT=cuda::std::__4::plus, OffsetT=int32_t, AccumT=int, StreamingContextT=cub::_V_300303_SM_1200::NullType]" at line 232 of /opt/compiler-explorer/cuda/13.3.0/bin/../targets/x86_64-linux/include/cccl/cub/device/dispatch/dispatch_reduce_by_key.cuh
[ 2 instantiation contexts not shown ]
instantiation of "cudaError_t cub::_V_300303_SM_1200::ChainedPolicy::find_and_invoke_policy(FunctorT &) [with PolicyPtxVersion=1000, PolicyT=cub::_V_300303_SM_1200::detail::rle::encode::policy_hub::Policy1000, PrevPolicyT=cub::_V_300303_SM_1200::detail::rle::encode::policy_hub::Policy900, DevicePtxVersion=1200, FunctorT=cub::_V_300303_SM_1200::detail::reduce::DispatchStreamingReduceByKey>, thrust::_V_300303_SM_1200::transform_output_iterator)->int>, int *>, thrust::_V_300303_SM_1200::constant_iterator, int *, int *, cuda::std::__4::equal_to, cuda::std::__4::plus, int32_t, int, cub::_V_300303_SM_1200::detail::rle::encode::policy_hub>]" at line 831 of /opt/compiler-explorer/cuda/13.3.0/bin/../targets/x86_64-linux/include/cccl/cub/util_device.cuh
instantiation of "cudaError_t cub::_V_300303_SM_1200::ChainedPolicy::runtime_arch_to_compiletime(int, FunctorT &) [with PolicyPtxVersion=1000, PolicyT=cub::_V_300303_SM_1200::detail::rle::encode::policy_hub::Policy1000, PrevPolicyT=cub::_V_300303_SM_1200::detail::rle::encode::policy_hub::Policy900, ArchMult=1, CudaArches=<1200>, FunctorT=cub::_V_300303_SM_1200::detail::reduce::DispatchStreamingReduceByKey>, thrust::_V_300303_SM_1200::transform_output_iterator)->int>, int *>, thrust::_V_300303_SM_1200::constant_iterator, int *, int *, cuda::std::__4::equal_to, cuda::std::__4::plus, int32_t, int, cub::_V_300303_SM_1200::detail::rle::encode::policy_hub>]" at line 796 of /opt/compiler-explorer/cuda/13.3.0/bin/../targets/x86_64-linux/include/cccl/cub/util_device.cuh
instantiation of "cudaError_t cub::_V_300303_SM_1200::ChainedPolicy::Invoke(int, FunctorT &) [with PolicyPtxVersion=1000, PolicyT=cub::_V_300303_SM_1200::detail::rle::encode::policy_hub::Policy1000, PrevPolicyT=cub::_V_300303_SM_1200::detail::rle::encode::policy_hub::Policy900, FunctorT=cub::_V_300303_SM_1200::detail::reduce::DispatchStreamingReduceByKey>, thrust::_V_300303_SM_1200::transform_output_iterator)->int>, int *>, thrust::_V_300303_SM_1200::constant_iterator, int *, int *, cuda::std::__4::equal_to, cuda::std::__4::plus, int32_t, int, cub::_V_300303_SM_1200::detail::rle::encode::policy_hub>]" at line 443 of /opt/compiler-explorer/cuda/13.3.0/bin/../targets/x86_64-linux/include/cccl/cub/device/dispatch/dispatch_streaming_reduce_by_key.cuh
instantiation of "cudaError_t cub::_V_300303_SM_1200::detail::reduce::DispatchStreamingReduceByKey::Dispatch(void *, size_t &, KeysInputIteratorT, UniqueOutputIteratorT, ValuesInputIteratorT, AggregatesOutputIteratorT, NumRunsOutputIteratorT, EqualityOpT, ReductionOpT, cub::_V_300303_SM_1200::detail::reduce::DispatchStreamingReduceByKey::global_offset_t, cudaStream_t) [with KeysInputIteratorT=thrust::_V_300303_SM_1200::zip_iterator>, UniqueOutputIteratorT=thrust::_V_300303_SM_1200::transform_output_iterator)->int>, int *>, ValuesInputIteratorT=thrust::_V_300303_SM_1200::constant_iterator, AggregatesOutputIteratorT=int *, NumRunsOutputIteratorT=int *, EqualityOpT=cuda::std::__4::equal_to, ReductionOpT=cuda::std::__4::plus, OffsetT=int32_t, AccumT=int, PolicyHub=cub::_V_300303_SM_1200::detail::rle::encode::policy_hub]" at line 191 of /opt/compiler-explorer/cuda/13.3.0/bin/../targets/x86_64-linux/include/cccl/cub/device/device_run_length_encode.cuh
instantiation of "cudaError_t cub::_V_300303_SM_1200::DeviceRunLengthEncode::Encode(void *, size_t &, InputIteratorT, UniqueOutputIteratorT, LengthsOutputIteratorT, NumRunsOutputIteratorT, NumItemsT, cudaStream_t) [with InputIteratorT=thrust::_V_300303_SM_1200::zip_iterator>, UniqueOutputIteratorT=thrust::_V_300303_SM_1200::transform_output_iterator)->int>, int *>, LengthsOutputIteratorT=int *, NumRunsOutputIteratorT=int *, NumItemsT=int]" at line 36 of
```
### How to Reproduce
https://godbolt.org/z/Prh5qTebT
```
#include
#include
#include
#include
#include
void func(){
int* items = nullptr;
int* segmentIds = nullptr;
auto rle_input = thrust::make_zip_iterator(
items,
segmentIds
);
int* unique_items = nullptr;
//does not compile
auto rle_output = thrust::make_transform_output_iterator(
unique_items,
cuda::proclaim_return_type(
[]__device__(cuda::std::tuple tup){
return cuda::std::get<0>(tup);
}
)
);
//does not compile
/*auto rle_output = thrust::make_zip_iterator(
unique_items,
thrust::make_discard_iterator()
);*/
size_t temp_storage_bytes = 0;
cub::DeviceRunLengthEncode::Encode(
nullptr,
temp_storage_bytes,
//(cuda::std::tuple*) nullptr,
rle_input,
//(cuda::std::tuple*) nullptr,
rle_output,
(int*)nullptr,
(int*)nullptr,
42
);
}
```
### Expected behavior
It works with transform_output_iterator
### Reproduction link
https://godbolt.org/z/Prh5qTebT
### Operating System
_No response_
### nvidia-smi output
_No response_
### NVCC version
_No response_
Contributor guide
Assessment
This issue has not been assessed yet.