cannot call std::sort with thrust::device_vector<T>::iterator

[andrewcorrigan]

This is related, but I think separate from #725. I cannot call std::sort using thrust::device_vector<T>::iterator. Compiling the below with AppleClang + libc++ I get the error message shown below, while compilation succeeds using AppleClang + libstdc++.

Reproducer:

#include <algorithm> 
#define THRUST_DEVICE_SYSTEM 4 // CPP 
#include <thrust/device_vector.h> 
int main(int argc, char** argv) 
{ 
    thrust::device_vector<int> a(10); 
    std::sort(a.begin(), a.end()); 
    return 0; 
}

Error message:

/Applications/Xcode.app/Contents/Developer/Toolchains/XcodeDefault.xctoolchain/usr/bin/../include/c++/v1/algorithm:3856:17: error: 
      no matching function for call to 'swap'
                swap(*__first, *__last);

where the relevant overload is

thrust/thrust/detail/device_reference.inl:49:6: note: 
      candidate function [with T = int] not viable: expects an l-value for 1st
      argument
void swap(device_reference<T> &a, device_reference<T> &b)

I tried adding overloads to thrust::swap, which instead accept either device_reference<T> or device_reference<T>&&, but then get the following error message. I'm just guessing at this point.

thrust/thrust/detail/reference.inl:196:3: error: 
      call to 'iter_swap' is ambiguous
  iter_swap(select_system(*system, *system), m_ptr, other.m_ptr); 

[jaredhoberock]

I think the signature of swap for device_reference may be insufficient, and it is leading to this issue.

Since swap(device_reference<T>& a, device_reference<T>& b) swaps the objects referred to by a and b, and not the pointers contained inside a and b, it's probably OK for swap() to receive the device_references by value.

Introducing another overload for free function swap which receives the device_references by value (not rvalue reference) ought to allow the compiler to find device_reference's swap() even when the arguments are temporaries.

I'm not sure what the error message regarding iter_swap implies.

[andrewcorrigan]

It looks like there are two implementations of iter_swap:

thrust/system/detail/generic/memory.inl:84:6: note: 
      candidate function [with Pointer1 = thrust::device_ptr<int>, Pointer2 =
      thrust::device_ptr<int>]
void iter_swap(tag, Pointer1, Pointer2)
     ^
thrust/system/detail/sequential/iter_swap.h:36:8: note: 
      candidate function [with Pointer1 = thrust::device_ptr<int>, Pointer2 =
      thrust::device_ptr<int>]
  void iter_swap(tag, Pointer1 a, Pointer2 b)

Only the one in sequential is implemented, so I tried qualifying the particular call to iter_swap in void reference<Element,Pointer,Derived>::swap(System *system, derived_type &other) as thrust::system::detail::sequential::iter_swap:

template<typename Element, typename Pointer, typename Derived>
  template<typename System>
    __host__ __device__
    void reference<Element,Pointer,Derived>
      ::swap(System *system, derived_type &other)
{
  using thrust::system::detail::generic::select_system;
  using thrust::system::detail::generic::iter_swap;

  thrust::system::detail::sequential::iter_swap(select_system(*system, *system), m_ptr, other.m_ptr);
} // end reference::swap()

and now the reproducer compiles.

I'd like to try to do a PR for this and also #725, both of which I need to get resolved for my code. Would you prefer swap(device_reference<T> a, device_reference<T> b) to be in addition to or instead of swap(device_reference<T>& a, device_reference<T>& b)?

[jaredhoberock]

reference::swap() needs to call iter_swap via ADL, so a workaround that calls it via qualification wouldn't be correct.

What backend are you compiling with?

[andrewcorrigan]

I compiled it so far with TBB and CPP.

[jaredhoberock]

I'm confused why the call would be ambiguous because I expect sequential::iter_swap() to be a closer match than generic::iter_swap.

[andrewcorrigan]

Is it because of the line above?

using thrust::system::detail::generic::iter_swap;

If instead of qualifying the call to iter_swap, I comment out using statement, then the reproducer also compiles.

[jaredhoberock]

That line is intended to be there. generic::iter_swap is only supposed to be chosen by ADL (and fail with a static assert) if no specialization of iter_swap exists. But a specialization does exist -- sequential::iter_swap.

[jaredhoberock]

I think the reason the call is ambiguous is that a conversion (to tag) has to happen to call either sequential::iter_swap() or generic::iter_swap(), and all conversions are considered equal. I don't recall why iter_swap()'s first parameter is a tag while the first parameter of other customization points are execution policies. My guess is that iter_swap() is assumed not to touch any state that may be stored inside an execution policy (but which a tag would discard).

The fact that this fails leads me to believe that swap(device_reference, device_reference) doesn't compile at all, independently of anything going on in sort().

Does a simple program that only tries to swap two device_references compile? If not, then we probably need to revisit the interface of iter_swap().

[andrewcorrigan]

[EDIT] This is wrong, see below.

Dereferencing a device_vector produces a device_reference, right? The following compiles, with either swap(device_reference<T> a, device_reference<T> b) or swap(device_reference<T>& a, device_reference<T>& b).

#define THRUST_DEVICE_SYSTEM 4    // CPP
#include <thrust/device_vector.h>
int main(int argc, char** argv)
{
    thrust::device_vector<int> x(10);
    auto a = *x.begin();
    auto b = *x.begin()+1;
    swap(a, b);
    return 0;
}

[jaredhoberock]

Then I really don't understand why your program above works, but the call to iter_swap() inside of reference::swap() fails. Your program above has to go through reference::swap().

[andrewcorrigan]

I messed up that example, one second

[andrewcorrigan]

b was defined incorrectly, sorry. This fails:

#include <algorithm>
#define THRUST_DEVICE_SYSTEM 4    // CPP
#include <thrust/device_vector.h>
int main(int argc, char** argv)
{
    thrust::device_vector<int> x(10);
    auto a = *x.begin();
    auto b = *(x.begin()+1);
    swap(a, b);
    return 0;
}

with the same iter_swap error:

thrust/detail/reference.inl:196:3: error: call to 'iter_swap' is ambiguous
  iter_swap(select_system(*system, *system), m_ptr, other.m_ptr);

[jaredhoberock]

OK, that makes more sense to me now. Does it only fail to compile with a particular compiler, or do all compilers agree that the program is illegal?

[andrewcorrigan]

I've been compiling with AppleClang from Xcode 8. g++-6 agrees that it's an error. icpc 17.0 also agrees.