[NVBug 3129879] MSVC debug iterators give dreaded "[subobject]" nvcc execution space warnings

[alliepiper]

Update:

This is an nvcc bug that will not be fixed. We'll need to disable the treat-warnings-as-errors flags for NVCC + MSVC CUDA Debug targets.

Original issue:

Encountered while trying to build thrust's bench.cu with MSVC 2019:

C:\PROGRA~1\NVIDIA~2\CUDA\v11.0\bin\nvcc.exe -forward-unknown-to-host-compiler -DNOMINMAX -DTHRUST_DEVICE_SYSTEM=THRUST_DEVICE_SYSTEM_CUDA -DTHRUST_HOST_SYSTEM=THRUST_HOST_SYSTEM_CPP -I..\ -I..\dependencies\cub -ftemplate-backtrace-limit 64  -gencode arch=compute_75,code=sm_75 -Xcompiler="-Zi -Ob0 -Od /RTC1" -Xcompiler=-MDd -Xcompiler=/W3 -Xcompiler=/WX -Xcompiler=/wd4244 -Xcompiler=/wd4267 -Xcompiler=/wd4800 -Xcompiler=/wd4146 -Xcompiler=/wd4494 -Xcompiler=/bigobj -Xcudafe=--display_error_number -Xcudafe=--promote_warnings -MD -MT internal\benchmark\CMakeFiles\thrust.cpp.cuda.cpp14.bench.dir\bench.cu.obj -MF internal\benchmark\CMakeFiles\thrust.cpp.cuda.cpp14.bench.dir\bench.cu.obj.d -x cu -c ..\internal\benchmark\bench.cu -o internal\benchmark\CMakeFiles\thrust.cpp.cuda.cpp14.bench.dir\bench.cu.obj -Xcompiler=-Fdinternal\benchmark\CMakeFiles\thrust.cpp.cuda.cpp14.bench.dir\,-FS
thrust/detail/allocator/allocator_traits.inl(396): error #3056-D: calling a __host__ function("std::_Iterator_base12::_Iterator_base12") from a __host__ __device__ function("std::_Vector_iterator<    ::std::_Vector_val<    ::std::_Simple_types<char> > > ::_Vector_iterator") is not allowed

Digging into this shows that the issue seems to be happening when std::vector iterators are passed to thrust::transform here and thrust::reduce here. I haven't yet tracked down what is causing the error.

This is similar to other issues, eg. NVIDIA/cccl#743. We need to make sure that we're actually testing all of the thrust:: algorithms with std::vector somewhere.

[alliepiper]

Branch that enables the benchmark code in CMake is here: https://github.com/allisonvacanti/thrust/tree/bug/add_bench_to_cmake/gh.1193

[alliepiper]

Just noticed this happening with other algorithms whenever doing an MSVC debug build. Seems related to the MSVC STL debug iterators.

[alliepiper]

This is unrelated to thrust:

repro.cu

#include <vector>

#pragma nv_exec_check_disable
template <typename T>
__host__ __device__
void hd_context()
{
  T t;
  (void)t;
}

int main()
{
  using cpu_iter_t = typename std::vector<int>::const_iterator;
  hd_context<cpu_iter_t>();
}

$ nvcc -g -Xcompiler -MDd repro.cu
scratch.cu
repro(10): warning: calling a __host__ function("std::_Iterator_base12::~_Iterator_base12") from a __host__ __device__ function("std::_Vector_const_iterator<    ::std::_Vector_val<    ::std::_Simple_types<int> > > ::~_Vector_const_iterator") is not allowed

repro(10): warning: calling a __host__ function("std::_Iterator_base12::~_Iterator_base12") from a __host__ __device__ function("std::_Vector_const_iterator<    ::std::_Vector_val<    ::std::_Simple_types<int> > > ::~_Vector_const_iterator [subobject]") is not allowed

   Creating library a.lib and object a.exp

This is on nvcc 11.0.167 and cl.exe 19.27.29111 (latest MSVC 2019).

[alliepiper]

Filed NVBug 3129879 requesting fix and/or workaround.

[alliepiper]

NVCC folks say that this is a limitation of the current implementation and cannot be fixed. The toolchain isn't equipped to handle MSVC's debug iterators.

Recommended workaround for now is to compile with --diag-suppress 20011, which is a new option in nvcc that will globally suppress all similar warnings.

Removing from current milestone since this is not a priority for us. Eventually we should update our build harness to just turn off -Werror all-warnings when we detect MSVC debug builds. We don't have many other options here.