compilation error when sorting bools with CUDA backend

[wnbell]

The problem is that radix sort is being dispatched, but the underlying radix sort implementation doesn't provide a bool specialization. Since we probably don't want to use radix sort here (it would be unstable for general numerical values in bool storage) we should dispatch partition or stable_partition instead.

Minimal reproducer:

#include <thrust/device_vector.h>
#include <thrust/sort.h>

int main(void)
{
  thrust::device_vector<bool> V(10);
  thrust::sort(V.begin(), V.end());
  return 0;
}

Compiler output:

/home/nathan/NV/thrust/thrust/system/cuda/detail/detail/stable_radix_sort.inl(67): error: incomplete type is not allowed
          detected during:
            instantiation of "void thrust::system::cuda::detail::detail::stable_radix_sort(RandomAccessIterator, RandomAccessIterator) [with RandomAccessIterator=thrust::detail::normal_iterator<thrust::pointer<__nv_bool, thrust::device_system_tag, thrust::use_default, thrust::use_default>>]" 
(62): here
            instantiation of "void thrust::system::cuda::detail::detail::stable_radix_sort(RandomAccessIterator, RandomAccessIterator) [with RandomAccessIterator=thrust::detail::normal_iterator<thrust::device_ptr<__nv_bool>>]" 
/home/nathan/NV/thrust/thrust/system/cuda/detail/sort.inl(284): here
            instantiation of "void thrust::system::cuda::detail::first_dispatch::stable_sort(RandomAccessIterator, RandomAccessIterator, StrictWeakOrdering, thrust::detail::true_type) [with RandomAccessIterator=thrust::detail::normal_iterator<thrust::device_ptr<__nv_bool>>, StrictWeakOrdering=thrust::less<__nv_bool>]" 
/home/nathan/NV/thrust/thrust/system/cuda/detail/sort.inl(419): here
            instantiation of "void thrust::system::cuda::detail::stable_sort(thrust::system::cuda::detail::tag, RandomAccessIterator, RandomAccessIterator, StrictWeakOrdering) [with RandomAccessIterator=thrust::detail::normal_iterator<thrust::device_ptr<__nv_bool>>, StrictWeakOrdering=thrust::less<__nv_bool>]" 
/home/nathan/NV/thrust/thrust/detail/sort.inl(88): here
            instantiation of "void thrust::stable_sort(RandomAccessIterator, RandomAccessIterator, StrictWeakOrdering) [with RandomAccessIterator=thrust::detail::normal_iterator<thrust::device_ptr<__nv_bool>>, StrictWeakOrdering=thrust::less<__nv_bool>]" 
/home/nathan/NV/thrust/thrust/system/detail/generic/sort.inl(59): here
            instantiation of "void thrust::system::detail::generic::sort(thrust::system::detail::generic::tag, RandomAccessIterator, RandomAccessIterator, StrictWeakOrdering) [with RandomAccessIterator=thrust::detail::normal_iterator<thrust::device_ptr<__nv_bool>>, StrictWeakOrdering=thrust::less<__nv_bool>]" 
/home/nathan/NV/thrust/thrust/detail/sort.inl(60): here
            instantiation of "void thrust::sort(RandomAccessIterator, RandomAccessIterator, StrictWeakOrdering) [with RandomAccessIterator=thrust::detail::normal_iterator<thrust::device_ptr<__nv_bool>>, StrictWeakOrdering=thrust::less<__nv_bool>]" 
/home/nathan/NV/thrust/thrust/system/detail/generic/sort.inl(47): here
            instantiation of "void thrust::system::detail::generic::sort(thrust::system::detail::generic::tag, RandomAccessIterator, RandomAccessIterator) [with RandomAccessIterator=thrust::detail::normal_iterator<thrust::device_ptr<__nv_bool>>]" 
/home/nathan/NV/thrust/thrust/detail/sort.inl(45): here
            instantiation of "void thrust::sort(RandomAccessIterator, RandomAccessIterator) [with RandomAccessIterator=thrust::detail::normal_iterator<thrust::device_ptr<__nv_bool>>]" 
foo.cu(7): here

/home/nathan/NV/thrust/thrust/system/cuda/detail/detail/stable_radix_sort.inl(67): error: incomplete type is not allowed
          detected during:
            instantiation of "void thrust::system::cuda::detail::detail::stable_radix_sort(RandomAccessIterator, RandomAccessIterator) [with RandomAccessIterator=thrust::detail::normal_iterator<thrust::device_ptr<__nv_bool>>]" 
/home/nathan/NV/thrust/thrust/system/cuda/detail/sort.inl(284): here
            instantiation of "void thrust::system::cuda::detail::first_dispatch::stable_sort(RandomAccessIterator, RandomAccessIterator, StrictWeakOrdering, thrust::detail::true_type) [with RandomAccessIterator=thrust::detail::normal_iterator<thrust::device_ptr<__nv_bool>>, StrictWeakOrdering=thrust::less<__nv_bool>]" 
/home/nathan/NV/thrust/thrust/system/cuda/detail/sort.inl(419): here
            instantiation of "void thrust::system::cuda::detail::stable_sort(thrust::system::cuda::detail::tag, RandomAccessIterator, RandomAccessIterator, StrictWeakOrdering) [with RandomAccessIterator=thrust::detail::normal_iterator<thrust::device_ptr<__nv_bool>>, StrictWeakOrdering=thrust::less<__nv_bool>]" 
/home/nathan/NV/thrust/thrust/detail/sort.inl(88): here
            instantiation of "void thrust::stable_sort(RandomAccessIterator, RandomAccessIterator, StrictWeakOrdering) [with RandomAccessIterator=thrust::detail::normal_iterator<thrust::device_ptr<__nv_bool>>, StrictWeakOrdering=thrust::less<__nv_bool>]" 
/home/nathan/NV/thrust/thrust/system/detail/generic/sort.inl(59): here
            instantiation of "void thrust::system::detail::generic::sort(thrust::system::detail::generic::tag, RandomAccessIterator, RandomAccessIterator, StrictWeakOrdering) [with RandomAccessIterator=thrust::detail::normal_iterator<thrust::device_ptr<__nv_bool>>, StrictWeakOrdering=thrust::less<__nv_bool>]" 
/home/nathan/NV/thrust/thrust/detail/sort.inl(60): here
            instantiation of "void thrust::sort(RandomAccessIterator, RandomAccessIterator, StrictWeakOrdering) [with RandomAccessIterator=thrust::detail::normal_iterator<thrust::device_ptr<__nv_bool>>, StrictWeakOrdering=thrust::less<__nv_bool>]" 
/home/nathan/NV/thrust/thrust/system/detail/generic/sort.inl(47): here
            instantiation of "void thrust::system::detail::generic::sort(thrust::system::detail::generic::tag, RandomAccessIterator, RandomAccessIterator) [with RandomAccessIterator=thrust::detail::normal_iterator<thrust::device_ptr<__nv_bool>>]" 
/home/nathan/NV/thrust/thrust/detail/sort.inl(45): here
            instantiation of "void thrust::sort(RandomAccessIterator, RandomAccessIterator) [with RandomAccessIterator=thrust::detail::normal_iterator<thrust::device_ptr<__nv_bool>>]" 
foo.cu(7): here

2 errors detected in the compilation of "/tmp/tmpxft_00005faf_00000000-4_foo.cpp1.ii".