How to use std::array in thrust::sort ?

[hzhangxyz]

See this code, After #3333 , I do not know how to use thrust::sort for std::array. However, in the former version, I can build it by nvcc main.cu -g -G -O0 -arch=sm_86 --expt-relaxed-constexpr, but now I cannot, and it reports error saying:

/opt/compiler-explorer/libs/cccl/trunk/cub/cub/thread/thread_sort.cuh(91): error: calling a __host__ function("std::enable_if<std::__array_traits<T1, T2> ::_Is_swappable::value, void> ::type  ::std::swap<int, (unsigned long)1ul> (    ::std::array<T1, T2>  &,     ::std::array<T1, T2>  &)") from a __device__ function("cub::CUB_300100_SM_700::StableOddEvenSort<    ::std::array<int, (unsigned long)1ul> ,  ::cub::CUB_300100_SM_700::NullType,     ::array_less<int, (long)1l> , (int)17> ") is not allowed

/opt/compiler-explorer/libs/cccl/trunk/cub/cub/thread/thread_sort.cuh(91): error: identifier "std::swap<int, (unsigned long)1ul> " is undefined in device code

[miscco]

The compile error you are receiving is correct.

std::array is a host only type and all accessor methods are host only. Consequently, thrust::device_vector<::std::array<...>> is an invalid type.

Luckily for you you can just use cuda::std::array from #include <cuda/std/array> as a drop in replacement.

[miscco]

A more elaborate answer would be that we need to swap elements around for sorting and std::swap which is found by ADL for std::array is not usable on device, so the compiler errors out.

[miscco]

Oh now I see the issue in the code.

Previously we would use cuda::std::swap for this, and that would be fine, because the compiler would generate the right swap function for std::array because it is a literal type.

However, that design is broken because any user calling "unqualified" swap, aka using std::swap; swap(a, b); will break if the type brings in both namespace std and namespace cuda::std.

For example std::array<cuda::std::complex> Here std::swap and cuda::std::swap would both be eligible for ADL and the program will not compile due to ambiguous function definitions.

[hzhangxyz]

I noticed that std::swap for std::array is constexpr since C++20, so if I set --std=c++20, this works!!!

[bernhardmgruber]

I think you were unfortunate to run into a problem due to our improvement. As @miscco said, the previous design, where we just called cub::Swap (or any other swap implementation we provided) was broken, because we did completely disregard any user provided swap(T, T) function for their type T. We fixed this by only using our own swap implementation if there is no other one available. This means that any type in namespace std::* will now correctly use std::swap. We are sorry that this came at a disadvantage for you!