NVIDIA / NVIDIA/cccl

[BUG]: Using cuda::std::zip_iterator with std::stable_sort runs into compilation error

Open
#6,400 1 comment 0 reactions 2 assignees Claimed by @oleksandr-pavlyk 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 bug and that I agree to the [Code of Conduct](CODE_OF_CONDUCT.md)

### Type of Bug

Compile-time Error

### Component

libcu++

### Describe the bug

Call to `std::stable_sort` with host-accessible `thrust::zip_iterator` works, but replacing it with `cuda::std::zip_iterator` runs into compilation errors.

### How to Reproduce

The reproducer is as follows:

```cu
#ifdef USE_CUDA_ITERATOR
# include
# include
# include
#else
# include
#endif

#include
#include

struct compare_first_lt_op_t
{
template
__host__ __device__ bool operator()(const LhsT& lhs, const RhsT& rhs) const
{
return thrust::get<0>(lhs) < thrust::get<0>(rhs);
}
};

int main(void) {
using key_t = int;
using data_t = long;

size_t num_items = 128;
thrust::host_vector keys_expected(num_items);
thrust::host_vector values_expected(num_items);

#ifdef USE_CUDA_ITERATOR
auto zipped_expected_it = cuda::make_zip_iterator(keys_expected.begin(), values_expected.begin());
#else
auto zipped_expected_it = thrust::make_zip_iterator(keys_expected.begin(), values_expected.begin());
#endif

std::stable_sort(zipped_expected_it, zipped_expected_it + num_items, compare_first_lt_op_t{});

return 0;
}
```

Compiling with `-UUSE_CUDA_ITERATOR` to use `thrust::zip_iterator` compiles fine. Switching to `-DUSE_CUDA_ITERATOR` to use `cuda::std::zip_iterator` runs into compilation errors.

Reproduced with CTK 12.9 and CTK 13.0.1, also with GCC 13 host compiler and LLVM 19 host compiler.

The compilation error log is below:

Compilation log

```
$ /usr/local/cuda/bin/nvcc test.cu --std=c++20 -arch=sm_86 -I ~/repos/cccl/thrust -I ~/repos/cccl/cub -I ~/repos/cccl/libcudacxx/include/ -DUSE_CUDA_ITERATOR
/usr/include/c++/13/bits/stl_algobase.h(185): error: no instance of overloaded function "std::swap" matches the argument list
argument types are: (cuda::std::__4::tuple, cuda::std::__4::tuple)
swap(*__a, *__b);
^
/usr/include/c++/13/bits/stl_pair.h(896): note #3327-D: candidate function template "std::swap(std::pair<_T1, _T2> &, std::pair<_T1, _T2> &)" failed deduction
swap(pair<_T1, _T2>&, pair<_T1, _T2>&) = delete;
^
/usr/include/c++/13/bits/stl_pair.h(879): note #3327-D: candidate function template "std::swap(std::pair<_T1, _T2> &, std::pair<_T1, _T2> &)" failed deduction
swap(pair<_T1, _T2>& __x, pair<_T1, _T2>& __y)
^
/usr/include/c++/13/bits/move.h(213): note #3327-D: candidate function template "std::swap(_Tp (&)[_Nm], _Tp (&)[_Nm])" failed deduction
swap(_Tp (&__a)[_Nm], _Tp (&__b)[_Nm])
^
/usr/include/c++/13/bits/move.h(189): note #3326-D: function template "std::swap(_Tp &, _Tp &)" does not match because argument #1 does not match parameter
swap(_Tp& __a, _Tp& __b)
^
/home/opavlyk/repos/cccl/libcudacxx/include/cuda/std/__memory/unique_ptr.h(516): note #3327-D: candidate function template "cuda::std::__4::swap(cuda::std::__4::unique_ptr<_Tp, _Dp> &, cuda::std::__4::unique_ptr<_Tp, _Dp> &) noexcept" failed deduction
swap(unique_ptr<_Tp, _Dp>& __x, unique_ptr<_Tp, _Dp>& __y) noexcept
^
/home/opavlyk/repos/cccl/libcudacxx/include/cuda/std/__memory/compressed_pair.h(215): note #3327-D: candidate function template "cuda::std::__4::swap(cuda::std::__4::__compressed_pair<_T1, _T2> &, cuda::std::__4::__compressed_pair<_T1, _T2> &)" failed deduction
__attribute__((__visibility__("hidden"))) constexpr void swap(__compressed_pair<_T1, _T2>& __x, __compressed_pair<_T1, _T2>& __y) noexcept(
^
/home/opavlyk/repos/cccl/libcudacxx/include/cuda/std/array(438): note #3327-D: candidate function template "cuda::std::__4::swap(cuda::std::__4::array<_Tp, _Size> &, cuda::std::__4::array<_Tp, _Size> &)" failed deduction
__attribute__((__visibility__("hidden"))) constexpr void swap(array<_Tp, _Size>& __x, array<_Tp, _Size>& __y) noexcept(noexcept(__x.swap(__y)))
^
/home/opavlyk/repos/cccl/libcudacxx/include/cuda/std/__optional/optional.h(849): note #3327-D: candidate function template "cuda::std::__4::swap(cuda::std::__4::optional<_Tp> &, cuda::std::__4::optional<_Tp> &)" failed deduction
swap(optional<_Tp>& __x, optional<_Tp>& __y) noexcept(noexcept(__x.swap(__y)))
^
/home/opavlyk/repos/cccl/libcudacxx/include/cuda/std/__utility/pair.h(613): note #3327-D: candidate function template "cuda::std::__4::swap(cuda::std::__4::pair<_T1, _T2> &, cuda::std::__4::pair<_T1, _T2> &)" failed deduction
swap(pair<_T1, _T2>& __x, pair<_T1, _T2>& __y) noexcept((is_nothrow_swappable_v<_T1> && is_nothrow_swappable_v<_T2>) )
^
/home/opavlyk/repos/cccl/libcudacxx/include/cuda/std/__tuple_dir/tuple_leaf.h(78): note #3327-D: candidate function template "cuda::std::__4::swap(cuda::std::__4::__tuple_leaf<_Ip, _Hp, _Ep> &, cuda::std::__4::__tuple_leaf<_Ip, _Hp, _Ep> &)" failed deduction
__attribute__((__visibility__("hidden"))) inline void swap(__tuple_leaf<_Ip, _Hp, _Ep>& __x,
^
/home/opavlyk/repos/cccl/libcudacxx/include/cuda/std/__utility/swap.h(52): note #3327-D: candidate function template "cuda::std::__4::swap(_Tp (&)[_Np], _Tp (&)[_Np])" failed deduction
swap(_Tp (&__a)[_Np], _Tp (&__b)[_Np]) noexcept(__is_nothrow_swappable<_Tp>::value)
^
/home/opavlyk/repos/cccl/libcudacxx/include/cuda/std/__utility/swap.h(41): note #3326-D: function template "cuda::std::__4::swap(_Tp &, cuda::std::__4::type_identity_t<_Tp> &)" does not match because argument #1 does not match parameter
__attribute__((__visibility__("hidden"))) constexpr __swap_result_t<_Tp> swap(_Tp& __x, type_identity_t<_Tp>& __y) noexcept(
^
detected during:
instantiation of "void std::iter_swap(_ForwardIterator1, _ForwardIterator2) [with _ForwardIterator1=cuda::__4::zip_iterator, thrust::THRUST_300200_SM_860_NS::detail::normal_iterator>, _ForwardIterator2=cuda::__4::zip_iterator, thrust::THRUST_300200_SM_860_NS::detail::normal_iterator>]" at line 2489 of /usr/include/c++/13/bits/stl_algo.h
instantiation of "void std::__merge_without_buffer(_BidirectionalIterator, _BidirectionalIterator, _BidirectionalIterator, _Distance, _Distance, _Compare) [with _BidirectionalIterator=cuda::__4::zip_iterator, thrust::THRUST_300200_SM_860_NS::detail::normal_iterator>, _Distance=long, _Compare=__gnu_cxx::__ops::_Iter_comp_iter]" at line 2799 of /usr/include/c++/13/bits/stl_algo.h
instantiation of "void std::__inplace_stable_sort(_RandomAccessIterator, _RandomAccessIterator, _Compare) [with _RandomAccessIterator=cuda::__4::zip_iterator, thrust::THRUST_300200_SM_860_NS::detail::normal_iterator>, _Compare=__gnu_cxx::__ops::_Iter_comp_iter]" at line 5045 of /usr/include/c++/13/bits/stl_algo.h
instantiation of "void std::__stable_sort(_RandomAccessIterator, _RandomAccessIterator, _Compare) [with _RandomAccessIterator=cuda::__4::zip_iterator, thrust::THRUST_300200_SM_860_NS::detail::normal_iterator>, _Compare=__gnu_cxx::__ops::_Iter_comp_iter]" at line 5120 of /usr/include/c++/13/bits/stl_algo.h
instantiation of "void std::stable_sort(_RAIter, _RAIter, _Compare) [with _RAIter=cuda::__4::zip_iterator, thrust::THRUST_300200_SM_860_NS::detail::normal_iterator>, _Compare=compare_first_lt_op_t]" at line 34 of test.cu

/usr/include/c++/13/bits/stl_algo.h(1399): error: no instance of overloaded function "std::__rotate" matches the argument list
argument types are: (cuda::__4::zip_iterator, thrust::THRUST_300200_SM_860_NS::detail::normal_iterator>, cuda::__4::zip_iterator, thrust::THRUST_300200_SM_860_NS::detail::normal_iterator>, cuda::__4::zip_iterator, thrust::THRUST_300200_SM_860_NS::detail::normal_iterator>, cuda::__4::__zv_iter_category_base_tag::iterator_category)
return std::__rotate(__first, __middle, __last,
^
/usr/include/c++/13/bits/stl_algo.h(1202): note #3326-D: function template "std::_V2::__rotate(_ForwardIterator, _ForwardIterator, _ForwardIterator, std::forward_iterator_tag)" does not match because argument #4 does not match parameter
__rotate(_ForwardIterator __first,
^
/usr/include/c++/13/bits/stl_algo.h(1244): note #3326-D: function template "std::_V2::__rotate(_BidirectionalIterator, _BidirectionalIterator, _BidirectionalIterator, std::bidirectional_iterator_tag)" does not match because argument #4 does not match parameter
__rotate(_BidirectionalIterator __first,
^
/usr/include/c++/13/bits/stl_algo.h(1283): note #3326-D: function template "std::_V2::__rotate(_RandomAccessIterator, _RandomAccessIterator, _RandomAccessIterator, std::random_access_iterator_tag)" does not match because argument #4 does not match parameter
__rotate(_RandomAccessIterator __first,
^
detected during:
instantiation of "_FIter std::_V2::rotate(_FIter, _FIter, _FIter) [with _FIter=cuda::__4::zip_iterator, thrust::THRUST_300200_SM_860_NS::detail::normal_iterator>]" at line 2517
instantiation of "void std::__merge_without_buffer(_BidirectionalIterator, _BidirectionalIterator, _BidirectionalIterator, _Distance, _Distance, _Compare) [with _BidirectionalIterator=cuda::__4::zip_iterator, thrust::THRUST_300200_SM_860_NS::detail::normal_iterator>, _Distance=long, _Compare=__gnu_cxx::__ops::_Iter_comp_iter]" at line 2799
instantiation of "void std::__inplace_stable_sort(_RandomAccessIterator, _RandomAccessIterator, _Compare) [with _RandomAccessIterator=cuda::__4::zip_iterator, thrust::THRUST_300200_SM_860_NS::detail::normal_iterator>, _Compare=__gnu_cxx::__ops::_Iter_comp_iter]" at line 5045
instantiation of "void std::__stable_sort(_RandomAccessIterator, _RandomAccessIterator, _Compare) [with _RandomAccessIterator=cuda::__4::zip_iterator, thrust::THRUST_300200_SM_860_NS::detail::normal_iterator>, _Compare=__gnu_cxx::__ops::_Iter_comp_iter]" at line 5120
instantiation of "void std::stable_sort(_RAIter, _RAIter, _Compare) [with _RAIter=cuda::__4::zip_iterator, thrust::THRUST_300200_SM_860_NS::detail::normal_iterator>, _Compare=compare_first_lt_op_t]" at line 34 of test.cu

2 errors detected in the compilation of "test.cu".
```

### Expected behavior

`cuda::std::make_zip_iterator` should work in place of `thrust::make_zip_iterator` for use in STL algorithms.

### Reproduction link

https://godbolt.org/z/8KxeWG19q

### Operating System

Ubuntu Linux 24.04

### nvidia-smi output

```
Thu Oct 30 10:10:27 2025
+-----------------------------------------------------------------------------------------+
| NVIDIA-SMI 575.57.08 Driver Version: 575.57.08 CUDA Version: 12.9 |
|-----------------------------------------+------------------------+----------------------+
| GPU Name Persistence-M | Bus-Id Disp.A | Volatile Uncorr. ECC |
| Fan Temp Perf Pwr:Usage/Cap | Memory-Usage | GPU-Util Compute M. |
| | | MIG M. |
|=========================================+========================+======================|
| 0 NVIDIA RTX A6000 On | 00000000:01:00.0 Off | Off |
| 31% 34C P8 18W / 300W | 1641MiB / 49140MiB | 29% Default |
| | | N/A |
+-----------------------------------------+------------------------+----------------------+
| 1 NVIDIA RTX A400 On | 00000000:C1:00.0 On | N/A |
| 30% 34C P8 N/A / 50W | 132MiB / 4094MiB | 0% Default |
| | | N/A |
+-----------------------------------------+------------------------+----------------------+
```

### NVCC version

```
$ /usr/local/cuda/bin/nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2025 NVIDIA Corporation
Built on Tue_May_27_02:21:03_PDT_2025
Cuda compilation tools, release 12.9, V12.9.86
Build cuda_12.9.r12.9/compiler.36037853_0
```

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.