Async sort algorithm test doesn't compile with MSVC

[alliepiper]

This looks to be a bug in cudafe, similar to nvbug 2732194:

../testing/async_sort.cu(148): error #871: template instantiation resulted in unexpected function type of 
  "thrust::event<thrust::cuda_cub::tag> (thrust::execution_policy<thrust::cuda_cub::tag> &,
                                         thrust::detail::normal_iterator<thrust::device_ptr<char>> &&,
                                         thrust::detail::normal_iterator<thrust::device_ptr<char>> &&, 
                                         thrust::less<char> &&)" 
    (the meaning of a name may have changed since the template declaration -- the type of the template is 
      "thrust::event<DerivedPolicy> (thrust::execution_policy<DerivedPolicy> &, 
                                     ForwardIt &&, 
                                     Sentinel &&, 
                                     StrictWeakOrdering &&)"
    )
          detected during:
            instantiation of "thrust::async::fallback::async_sort" based on template arguments <thrust::cuda_cub::tag, thrust::detail::normal_iterator<thrust::device_ptr<char>>, thrust::detail::normal_iterator<thrust::device_ptr<char>>, thrust::less<char>> 
(148): here
            instantiation of "void test_async_sort<SortInvoker, WaitPolicy>::tester<T>::operator()(size_t) [with SortInvoker=sort_invoker, WaitPolicy=wait_for_futures, T=char]" 
../testing\unittest/meta.h(104): here
            instantiation of "void unittest::for_each_type<TypeList, Function, T, i>::operator()(U) [with TypeList=NumericTypes, Function=test_async_sort<sort_invoker, wait_for_futures>::tester, T=char, i=0U, U=size_t]" 
../testing\unittest/testframework.h(534): here
            instantiation of "void VariableUnitTest<TestName, TypeList>::run() [with TestName=test_async_sort<sort_invoker, wait_for_futures>::tester, TypeList=NumericTypes]" 

This is triaged by disabling the async copy algorithm + test on MSVC, but needs further investigation.

[RAMitchell]

I am having similar issues with async reduction when attempting to migrate our xgboost codebase to new asynchronous APIs. Any compile failures on MSVC will be problematic for us as we will be stuck with performance regressions from old synchronous functions.

[alliepiper]

@RAMitchell Are you seeing the template instantiation resulted in unexpected function type of ... error in the async reduce algorithm? If you have a small example that reproduces it, I'll take a look.

As a heads up, the async algorithms haven't gotten a ton of testing on MSVC. We're working on improving CI to shake out bugs, but those algorithms on that compiler might be flakey for a little while.

[RAMitchell]

If I compile this

#include <thrust/async/reduce.h>
#include <thrust/device_vector.h>

int main()
{
  int n = 5;
  thrust::device_vector<int> x(n, 1);
  auto result = thrust::async::reduce(x.begin(), x.end());
}

Adding the following as a compile argument to convince thrust that c++11 exists:
-Xcompiler "/Zc:__cplusplus"

I get the following:

1>------ Build started: Project: thrust_async, Configuration: Debug x64 ------
1>Compiling CUDA source file kernel.cu...
1>
1>C:\Users\Rory\source\repos\thrust_async\thrust_async>"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v10.2\bin\nvcc.exe" -gencode=arch=compute_30,code=\"sm_30,compute_30\" --use-local-env -ccbin "C:\Program Files (x86)\Microsoft Visual Studio\2017\Community\VC\Tools\MSVC\14.16.27023\bin\HostX86\x64" -x cu  -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v10.2\include" -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v10.2\include"  -G   --keep-dir x64\Debug -maxrregcount=0  --machine 64 --compile -cudart static -Xcompiler "/Zc:__cplusplus" -g   -DWIN32 -DWIN64 -D_DEBUG -D_CONSOLE -D_MBCS -Xcompiler "/EHsc /W3 /nologo /Od /Fdx64\Debug\vc141.pdb /FS /Zi /RTC1 /MDd " -o x64\Debug\kernel.cu.obj "C:\Users\Rory\source\repos\thrust_async\thrust_async\kernel.cu"
1>C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v10.2\include\thrust/detail/execute_with_dependencies.h(193): warning : __declspec attributes ignored
1>
1>C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v10.2\include\thrust/detail/execute_with_dependencies.h(193): warning : __declspec attributes ignored
1>
1>C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v10.2\include\thrust/system/cuda/detail/async/reduce.h(77): error : template parameter "thrust::detail::allocator_traits<Alloc>::rebind_traits [with Alloc=const auto]" may not have a template argument list
1>
1>1 error detected in the compilation of "C:/Users/Rory/AppData/Local/Temp/tmpxft_00004884_00000000-5_kernel.cpp4.ii".
1>kernel.cu
1>C:\Program Files (x86)\Microsoft Visual Studio\2017\Community\Common7\IDE\VC\VCTargets\BuildCustomizations\CUDA 10.2.targets(764,9): error MSB3721: The command ""C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v10.2\bin\nvcc.exe" -gencode=arch=compute_30,code=\"sm_30,compute_30\" --use-local-env -ccbin "C:\Program Files (x86)\Microsoft Visual Studio\2017\Community\VC\Tools\MSVC\14.16.27023\bin\HostX86\x64" -x cu  -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v10.2\include" -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v10.2\include"  -G   --keep-dir x64\Debug -maxrregcount=0  --machine 64 --compile -cudart static -Xcompiler "/Zc:__cplusplus" -g   -DWIN32 -DWIN64 -D_DEBUG -D_CONSOLE -D_MBCS -Xcompiler "/EHsc /W3 /nologo /Od /Fdx64\Debug\vc141.pdb /FS /Zi /RTC1 /MDd " -o x64\Debug\kernel.cu.obj "C:\Users\Rory\source\repos\thrust_async\thrust_async\kernel.cu"" exited with code 2.
1>Done building project "thrust_async.vcxproj" -- FAILED.
========== Build: 0 succeeded, 1 failed, 0 up-to-date, 0 skipped ==========

This is a problem for us, I have to migrate the entire xgboost code base off thrust onto cub until synchronisation issues are resolved. See dmlc/xgboost#5560.

[alliepiper]

Ah, gotcha, that's a different bug than the one in this github issue. That compile error (and many other issues on MSVC) are fixed in #1089 and thrust/cub#21, which I hope to have merged soon. If you get a chance to try those out, let me know if you find any other issues.

[alliepiper]

Also, -Xcompiler "/Zc:__cplusplus" is no longer necessary on that branch. We will properly detect the C++ standard now on supported versions of MSVC.

[alliepiper]

The issue in the description can be reproduced with the following minimal reducer on:

• MSVC 2017 x64
• CUDA 10.2

**********************************************************************
** Visual Studio 2017 Developer Command Prompt v15.9.21
** Copyright (c) 2017 Microsoft Corporation
**********************************************************************
[vcvarsall.bat] Environment initialized for: 'x64'

C:\Program Files (x86)\Microsoft Visual Studio\2017\Community>cd \Users\Allie\src\repro-thrust-1098

C:\Users\allie\src\repro-thrust-1098>"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v10.2\bin\nvcc.exe" -c repro.cu
repro.cu
repro.cu(83): error: template instantiation resulted in unexpected function type of "thrust::event<thrust::cuda::par_t> (thrust::cuda::par_t)" (the meaning of a name may have changed since the template declaration -- the type of the template is "thrust::event<T> (T)")
          detected during instantiation of "thrust::fallback::async_sort" based on template argument <thrust::cuda::par_t>
(83): here

repro.cu(83): error: no instance of function template "thrust::sort_detail::sort_fn::call" matches the argument list
            argument types are: (const thrust::cuda::par_t)

2 errors detected in the compilation of "C:/Users/allie/AppData/Local/Temp/tmpxft_00005444_00000000-10_repro.cpp1.ii".

repro.cu:

#include <utility> // for declval

namespace thrust
{

template <typename DerivedPolicy>
struct execution_policy {};

namespace cuda
{

template <class Derived>
struct execution_policy : thrust::execution_policy<Derived> {};

struct par_t : execution_policy<par_t> {};

static constexpr par_t par;

struct unique_eager_event {};

} // namespace cuda

// Unimplemented, only used in decltype contexts
template <typename DerivedPolicy>
thrust::cuda::unique_eager_event
unique_eager_event_type(thrust::cuda::execution_policy<DerivedPolicy> const&) noexcept;

namespace unimplemented
{

struct no_unique_eager_event_type_found {};

inline __host__
no_unique_eager_event_type_found
unique_eager_event_type(...) noexcept;

} // namespace unimplemented

namespace unique_eager_event_type_detail
{

using unimplemented::unique_eager_event_type;

template <typename System>
using select = decltype(unique_eager_event_type(std::declval<System>()));

} // namespace unique_eager_event_type_detail

template <typename System>
using unique_eager_event = unique_eager_event_type_detail::select<System>;

template <typename System>
using event = unique_eager_event<System>;

namespace fallback
{

template <typename T>
thrust::event<T> async_sort(T val) { return {}; }

} // namespace fallback

namespace sort_detail
{

using fallback::async_sort;

struct sort_fn final
{
  template <typename T>
  static auto call(T exec) -> decltype(async_sort(exec))
  {
    return async_sort(exec);
  }
};

} // namespace sort_detail

} // end namespace thrust

void foo()
{
  thrust::sort_detail::sort_fn::call(thrust::cuda::par);
}

[alliepiper]

Filed an internal bug against nvcc: nvbug id 2947213

[alliepiper]

Just confirmed that this will be fixed in a future version of NVCC.

Need to update the guard at the top of async_sort.cu to only disable the test on nvcc && nvcc < (tentatively) 11.1.

[alliepiper]

#1225 re-enables this test on MSVC when nvcc > 11.0 detected.