Can't build with NVCC option '--Werror cross-execution-space-call' on Windows

[Robadob]

Tested under Visual Studio 2015 and 2019, with CUDA 10.1/10.2.

If you try to build the below example with --Werror cross-execution-space-call enabled in Visual Studio it will fail.

This is the flag for error : calling a __host__ function from a __host__ __device__ function is not allowed

#include "jitify/jitify.hpp"

int main() {
    return EXIT_SUCCESS;
}

Under Visual Studio 2015, it produces the compilation error shown at the bottom of this issue.
TLDR: This is the offending line _ranked_keys.erase(rank);

I'm not actually sure why it's being triggered in this instance, but it's upsetting one of our continuous integration builds, so it would nice if it could be fixed without having to disable that compiler argument

1>  C:\Users\rob\Documents\Visual Studio 2015\Projects\Jitifytest>"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v10.2\bin\nvcc.exe" -gencode=arch=compute_61,code=\"sm_61,compute_61\" --use-local-env -ccbin "C:\Program Files (x86)\Microsoft Visual Studio 14.0\VC\bin\x86_amd64" -x cu -rdc=true -I"C:\Program Files (x86)\dlfcn-win32\include" -I"C:\Program Files (x86)\Visual Leak Detector\include" -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 --Werror cross-execution-space-call -g   -DWIN32 -DWIN64 -D_DEBUG -D_CONSOLE -D_MBCS -Xcompiler "/EHsc /W3 /nologo /Od /Fdx64\Debug\vc140.pdb /FS /Zi /RTC1 /MDd " -o x64\Debug\kernel.cu.obj "C:\Users\rob\Documents\Visual Studio 2015\Projects\Jitifytest\kernel.cu"
1>C:\Program Files (x86)\Microsoft Visual Studio 14.0\VC\include\xutility(2636): error : calling a __host__ function from a __host__ __device__ function is not allowed
1>            detected during:
1>              instantiation of "_OutIt std::move(_InIt, _InIt, _OutIt) [with _InIt=std::_Deque_iterator<std::_Deque_val<std::_Deque_simple_types<jitify::JitCache_impl::key_type>>>, _OutIt=std::_Deque_iterator<std::_Deque_val<std::_Deque_simple_types<jitify::JitCache_impl::key_type>>>]"
1>  C:\Program Files (x86)\Microsoft Visual Studio 14.0\VC\include\deque(1660): here
1>              instantiation of "std::deque<_Ty, _Alloc>::iterator std::deque<_Ty, _Alloc>::erase(std::deque<_Ty, _Alloc>::const_iterator, std::deque<_Ty, _Alloc>::const_iterator) [with _Ty=jitify::JitCache_impl::key_type, _Alloc=std::allocator<jitify::JitCache_impl::key_type>]"
1>  C:\Program Files (x86)\Microsoft Visual Studio 14.0\VC\include\deque(1628): here
1>              instantiation of "std::deque<_Ty, _Alloc>::iterator std::deque<_Ty, _Alloc>::erase(std::deque<_Ty, _Alloc>::const_iterator) [with _Ty=jitify::JitCache_impl::key_type, _Alloc=std::allocator<jitify::JitCache_impl::key_type>]"
1>  c:\users\rob\documents\visual studio 2015\projects\jitifytest\jitify/jitify.hpp(246): here
1>              instantiation of "void jitify::ObjectCache<KeyType, ValueType>::touch(const jitify::ObjectCache<KeyType, ValueType>::key_type &) [with KeyType=jitify::JitCache_impl::key_type, ValueType=jitify::detail::CUDAKernel]"
1>  c:\users\rob\documents\visual studio 2015\projects\jitifytest\jitify/jitify.hpp(254): here
1>              instantiation of "jitify::ObjectCache<KeyType, ValueType>::value_type &jitify::ObjectCache<KeyType, ValueType>::get(const jitify::ObjectCache<KeyType, ValueType>::key_type &) [with KeyType=jitify::JitCache_impl::key_type, ValueType=jitify::detail::CUDAKernel]"
1>  c:\users\rob\documents\visual studio 2015\projects\jitifytest\jitify/jitify.hpp(2861): here
1>
1>C:\Program Files (x86)\Microsoft Visual Studio 14.0\VC\include\xutility(2636): error : calling a __host__ function from a __host__ __device__ function is not allowed
1>            detected during:
1>              instantiation of "_OutIt std::move(_InIt, _InIt, _OutIt) [with _InIt=std::_Deque_iterator<std::_Deque_val<std::_Deque_simple_types<jitify::JitCache_impl::key_type>>>, _OutIt=std::_Deque_iterator<std::_Deque_val<std::_Deque_simple_types<jitify::JitCache_impl::key_type>>>]"
1>  C:\Program Files (x86)\Microsoft Visual Studio 14.0\VC\include\deque(1660): here
1>              instantiation of "std::deque<_Ty, _Alloc>::iterator std::deque<_Ty, _Alloc>::erase(std::deque<_Ty, _Alloc>::const_iterator, std::deque<_Ty, _Alloc>::const_iterator) [with _Ty=jitify::JitCache_impl::key_type, _Alloc=std::allocator<jitify::JitCache_impl::key_type>]"
1>  C:\Program Files (x86)\Microsoft Visual Studio 14.0\VC\include\deque(1628): here
1>              instantiation of "std::deque<_Ty, _Alloc>::iterator std::deque<_Ty, _Alloc>::erase(std::deque<_Ty, _Alloc>::const_iterator) [with _Ty=jitify::JitCache_impl::key_type, _Alloc=std::allocator<jitify::JitCache_impl::key_type>]"
1>  c:\users\rob\documents\visual studio 2015\projects\jitifytest\jitify/jitify.hpp(246): here
1>              instantiation of "void jitify::ObjectCache<KeyType, ValueType>::touch(const jitify::ObjectCache<KeyType, ValueType>::key_type &) [with KeyType=jitify::JitCache_impl::key_type, ValueType=jitify::detail::CUDAKernel]"
1>  c:\users\rob\documents\visual studio 2015\projects\jitifytest\jitify/jitify.hpp(254): here
1>              instantiation of "jitify::ObjectCache<KeyType, ValueType>::value_type &jitify::ObjectCache<KeyType, ValueType>::get(const jitify::ObjectCache<KeyType, ValueType>::key_type &) [with KeyType=jitify::JitCache_impl::key_type, ValueType=jitify::detail::CUDAKernel]"
1>  c:\users\rob\documents\visual studio 2015\projects\jitifytest\jitify/jitify.hpp(2861): here
1>
1>C:\Program Files (x86)\Microsoft Visual Studio 14.0\VC\include\xutility(2637): error : calling a __host__ function from a __host__ __device__ function is not allowed
1>            detected during:
1>              instantiation of "_OutIt std::move(_InIt, _InIt, _OutIt) [with _InIt=std::_Deque_iterator<std::_Deque_val<std::_Deque_simple_types<jitify::JitCache_impl::key_type>>>, _OutIt=std::_Deque_iterator<std::_Deque_val<std::_Deque_simple_types<jitify::JitCache_impl::key_type>>>]"
1>  C:\Program Files (x86)\Microsoft Visual Studio 14.0\VC\include\deque(1660): here
1>              instantiation of "std::deque<_Ty, _Alloc>::iterator std::deque<_Ty, _Alloc>::erase(std::deque<_Ty, _Alloc>::const_iterator, std::deque<_Ty, _Alloc>::const_iterator) [with _Ty=jitify::JitCache_impl::key_type, _Alloc=std::allocator<jitify::JitCache_impl::key_type>]"
1>  C:\Program Files (x86)\Microsoft Visual Studio 14.0\VC\include\deque(1628): here
1>              instantiation of "std::deque<_Ty, _Alloc>::iterator std::deque<_Ty, _Alloc>::erase(std::deque<_Ty, _Alloc>::const_iterator) [with _Ty=jitify::JitCache_impl::key_type, _Alloc=std::allocator<jitify::JitCache_impl::key_type>]"
1>  c:\users\rob\documents\visual studio 2015\projects\jitifytest\jitify/jitify.hpp(246): here
1>              instantiation of "void jitify::ObjectCache<KeyType, ValueType>::touch(const jitify::ObjectCache<KeyType, ValueType>::key_type &) [with KeyType=jitify::JitCache_impl::key_type, ValueType=jitify::detail::CUDAKernel]"
1>  c:\users\rob\documents\visual studio 2015\projects\jitifytest\jitify/jitify.hpp(254): here
1>              instantiation of "jitify::ObjectCache<KeyType, ValueType>::value_type &jitify::ObjectCache<KeyType, ValueType>::get(const jitify::ObjectCache<KeyType, ValueType>::key_type &) [with KeyType=jitify::JitCache_impl::key_type, ValueType=jitify::detail::CUDAKernel]"
1>  c:\users\rob\documents\visual studio 2015\projects\jitifytest\jitify/jitify.hpp(2861): here
1>
1>  3 errors detected in the compilation of "C:/Users/rob/AppData/Local/Temp/tmpxft_000033b4_00000000-7_kernel.cpp1.ii".
1>  kernel.cu

(Note jitify.hpp(2861), shown at the bottom is an incorrect line number, the length of jitify.hpp confuses visual studio compiler for some reason, i previously looked into this and couldn't find an explicit cause beyond breaking the file up into smaller files corrects the line numbers.)

[mondus]

There is also an issue with assignment in a conditional statement. I have issued a PR (#63) for that but am not sure how to fix the issue above.

[benbarsdell]

Thanks for the PR!

I filed an internal bug about the __host__ __device__ warnings; it seems to be a compiler issue.
I believe it only affects debug builds, but I don't know if that's of much help to you.

[ptheywood]

The __host__ __device__ warnings through MSVC also effect our release builds:

"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v10.1\bin\nvcc.exe" -gencode=arch=compute_35,code=\"sm_35,compute_35\" -gencode=arch=compute_60,code=\"sm_60,compute_60\" -gencode=arch=compute_60,code=\"compute_60,compute_60\" --use-local-env -ccbin "C:\Program Files (x86)\Microsoft Visual Studio\2019\Community\VC\Tools\MSVC\14.25.28610\bin\HostX64\x64" -x cu -rdc=true -IC:\projects\FLAMEGPU2_dev\include -IC:\projects\FLAMEGPU2_dev\src -IC:\projects\FLAMEGPU2_dev\externals -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v10.1\include\..\include" -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v10.1\include" -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v10.1\include"     --keep-dir x64\Release -maxrregcount=0  --machine 64 --compile -cudart static --expt-relaxed-constexpr --Wreorder --Werror reorder,cross-execution-space-call -Xptxas=-Werror -Xnvlink=-Werror -lineinfo -Xcompiler="/EHsc /wd4505 /WX -Ob2"    -D_WINDOWS -DJITIFY_PRINT_LOG -DNDEBUG -D"CMAKE_INTDIR=\"Release\"" -DWIN32 -D_WINDOWS -DJITIFY_PRINT_LOG -DNDEBUG -D"CMAKE_INTDIR=\"Release\"" -D_MBCS -Xcompiler "/EHsc /W4 /nologo /O2 /Fdflamegpu2.dir\Release\flamegpu2.pdb /FS /Zi  /MD /GR" -o flamegpu2.dir\Release\CUDAScanCompaction.obj "C:\projects\FLAMEGPU2_dev\src\flamegpu\gpu\CUDAScanCompaction.cu" 
C:\Program Files (x86)\Microsoft Visual Studio\2019\Community\VC\Tools\MSVC\14.25.28610\include\xutility(3942): error : calling a __host__ function from a __host__ __device__ function is not allowed [C:\projects\FLAMEGPU2_dev\build\FLAMEGPU2\flamegpu2.vcxproj]
            detected during:
              instantiation of "_OutIt std::move(_InIt, _InIt, _OutIt) [with _InIt=std::_Deque_iterator<std::_Deque_val<std::conditional_t<true, std::_Deque_simple_types<jitify::JitCache_impl::key_type>, std::_Deque_iter_types<jitify::JitCache_impl::key_type, size_t, ptrdiff_t, jitify::JitCache_impl::key_type *, const jitify::JitCache_impl::key_type *, jitify::JitCache_impl::key_type &, const jitify::JitCache_impl::key_type &, jitify::JitCache_impl::key_type **>>>>, _OutIt=std::_Deque_iterator<std::_Deque_val<std::conditional_t<true, std::_Deque_simple_types<jitify::JitCache_impl::key_type>, std::_Deque_iter_types<jitify::JitCache_impl::key_type, size_t, ptrdiff_t, jitify::JitCache_impl::key_type *, const jitify::JitCache_impl::key_type *, jitify::JitCache_impl::key_type &, const jitify::JitCache_impl::key_type &, jitify::JitCache_impl::key_type **>>>>]" 
  C:\Program Files (x86)\Microsoft Visual Studio\2019\Community\VC\Tools\MSVC\14.25.28610\include\deque(1302): here
              instantiation of "std::deque<_Ty, _Alloc>::iterator std::deque<_Ty, _Alloc>::erase(std::deque<_Ty, _Alloc>::const_iterator, std::deque<_Ty, _Alloc>::const_iterator) [with _Ty=jitify::JitCache_impl::key_type, _Alloc=std::allocator<jitify::JitCache_impl::key_type>]" 
  C:\Program Files (x86)\Microsoft Visual Studio\2019\Community\VC\Tools\MSVC\14.25.28610\include\deque(1275): here
              instantiation of "std::deque<_Ty, _Alloc>::iterator std::deque<_Ty, _Alloc>::erase(std::deque<_Ty, _Alloc>::const_iterator) [with _Ty=jitify::JitCache_impl::key_type, _Alloc=std::allocator<jitify::JitCache_impl::key_type>]" 
  C:\projects\FLAMEGPU2_dev\externals\jitify/jitify.hpp(227): here
              instantiation of "void jitify::ObjectCache<KeyType, ValueType>::touch(const jitify::ObjectCache<KeyType, ValueType>::key_type &) [with KeyType=jitify::JitCache_impl::key_type, ValueType=jitify::detail::CUDAKernel]" 
  C:\projects\FLAMEGPU2_dev\externals\jitify/jitify.hpp(235): here
              instantiation of "jitify::ObjectCache<KeyType, ValueType>::value_type &jitify::ObjectCache<KeyType, ValueType>::get(const jitify::ObjectCache<KeyType, ValueType>::key_type &) [with KeyType=jitify::JitCache_impl::key_type, ValueType=jitify::detail::CUDAKernel]" 
  C:\projects\FLAMEGPU2_dev\externals\jitify/jitify.hpp(2866): here
  
C:\Program Files (x86)\Microsoft Visual Studio\2019\Community\VC\Tools\MSVC\14.25.28610\include\xutility(3943): error : calling a __host__ function from a __host__ __device__ function is not allowed [C:\projects\FLAMEGPU2_dev\build\FLAMEGPU2\flamegpu2.vcxproj]
            detected during:
              instantiation of "_OutIt std::move(_InIt, _InIt, _OutIt) [with _InIt=std::_Deque_iterator<std::_Deque_val<std::conditional_t<true, std::_Deque_simple_types<jitify::JitCache_impl::key_type>, std::_Deque_iter_types<jitify::JitCache_impl::key_type, size_t, ptrdiff_t, jitify::JitCache_impl::key_type *, const jitify::JitCache_impl::key_type *, jitify::JitCache_impl::key_type &, const jitify::JitCache_impl::key_type &, jitify::JitCache_impl::key_type **>>>>, _OutIt=std::_Deque_iterator<std::_Deque_val<std::conditional_t<true, std::_Deque_simple_types<jitify::JitCache_impl::key_type>, std::_Deque_iter_types<jitify::JitCache_impl::key_type, size_t, ptrdiff_t, jitify::JitCache_impl::key_type *, const jitify::JitCache_impl::key_type *, jitify::JitCache_impl::key_type &, const jitify::JitCache_impl::key_type &, jitify::JitCache_impl::key_type **>>>>]" 
  C:\Program Files (x86)\Microsoft Visual Studio\2019\Community\VC\Tools\MSVC\14.25.28610\include\deque(1302): here
              instantiation of "std::deque<_Ty, _Alloc>::iterator std::deque<_Ty, _Alloc>::erase(std::deque<_Ty, _Alloc>::const_iterator, std::deque<_Ty, _Alloc>::const_iterator) [with _Ty=jitify::JitCache_impl::key_type, _Alloc=std::allocator<jitify::JitCache_impl::key_type>]" 
  C:\Program Files (x86)\Microsoft Visual Studio\2019\Community\VC\Tools\MSVC\14.25.28610\include\deque(1275): here
              instantiation of "std::deque<_Ty, _Alloc>::iterator std::deque<_Ty, _Alloc>::erase(std::deque<_Ty, _Alloc>::const_iterator) [with _Ty=jitify::JitCache_impl::key_type, _Alloc=std::allocator<jitify::JitCache_impl::key_type>]" 
  C:\projects\FLAMEGPU2_dev\externals\jitify/jitify.hpp(227): here
              instantiation of "void jitify::ObjectCache<KeyType, ValueType>::touch(const jitify::ObjectCache<KeyType, ValueType>::key_type &) [with KeyType=jitify::JitCache_impl::key_type, ValueType=jitify::detail::CUDAKernel]" 
  C:\projects\FLAMEGPU2_dev\externals\jitify/jitify.hpp(235): here
              instantiation of "jitify::ObjectCache<KeyType, ValueType>::value_type &jitify::ObjectCache<KeyType, ValueType>::get(const jitify::ObjectCache<KeyType, ValueType>::key_type &) [with KeyType=jitify::JitCache_impl::key_type, ValueType=jitify::detail::CUDAKernel]" 
  C:\projects\FLAMEGPU2_dev\externals\jitify/jitify.hpp(2866): here
  
C:\Program Files (x86)\Microsoft Visual Studio\2019\Community\VC\Tools\MSVC\14.25.28610\include\xutility(3944): error : calling a __host__ function from a __host__ __device__ function is not allowed [C:\projects\FLAMEGPU2_dev\build\FLAMEGPU2\flamegpu2.vcxproj]
            detected during:
              instantiation of "_OutIt std::move(_InIt, _InIt, _OutIt) [with _InIt=std::_Deque_iterator<std::_Deque_val<std::conditional_t<true, std::_Deque_simple_types<jitify::JitCache_impl::key_type>, std::_Deque_iter_types<jitify::JitCache_impl::key_type, size_t, ptrdiff_t, jitify::JitCache_impl::key_type *, const jitify::JitCache_impl::key_type *, jitify::JitCache_impl::key_type &, const jitify::JitCache_impl::key_type &, jitify::JitCache_impl::key_type **>>>>, _OutIt=std::_Deque_iterator<std::_Deque_val<std::conditional_t<true, std::_Deque_simple_types<jitify::JitCache_impl::key_type>, std::_Deque_iter_types<jitify::JitCache_impl::key_type, size_t, ptrdiff_t, jitify::JitCache_impl::key_type *, const jitify::JitCache_impl::key_type *, jitify::JitCache_impl::key_type &, const jitify::JitCache_impl::key_type &, jitify::JitCache_impl::key_type **>>>>]" 
  C:\Program Files (x86)\Microsoft Visual Studio\2019\Community\VC\Tools\MSVC\14.25.28610\include\deque(1302): here
              instantiation of "std::deque<_Ty, _Alloc>::iterator std::deque<_Ty, _Alloc>::erase(std::deque<_Ty, _Alloc>::const_iterator, std::deque<_Ty, _Alloc>::const_iterator) [with _Ty=jitify::JitCache_impl::key_type, _Alloc=std::allocator<jitify::JitCache_impl::key_type>]" 
  C:\Program Files (x86)\Microsoft Visual Studio\2019\Community\VC\Tools\MSVC\14.25.28610\include\deque(1275): here
              instantiation of "std::deque<_Ty, _Alloc>::iterator std::deque<_Ty, _Alloc>::erase(std::deque<_Ty, _Alloc>::const_iterator) [with _Ty=jitify::JitCache_impl::key_type, _Alloc=std::allocator<jitify::JitCache_impl::key_type>]" 
  C:\projects\FLAMEGPU2_dev\externals\jitify/jitify.hpp(227): here
              instantiation of "void jitify::ObjectCache<KeyType, ValueType>::touch(const jitify::ObjectCache<KeyType, ValueType>::key_type &) [with KeyType=jitify::JitCache_impl::key_type, ValueType=jitify::detail::CUDAKernel]" 
  C:\projects\FLAMEGPU2_dev\externals\jitify/jitify.hpp(235): here
              instantiation of "jitify::ObjectCache<KeyType, ValueType>::value_type &jitify::ObjectCache<KeyType, ValueType>::get(const jitify::ObjectCache<KeyType, ValueType>::key_type &) [with KeyType=jitify::JitCache_impl::key_type, ValueType=jitify::detail::CUDAKernel]" 
  C:\projects\FLAMEGPU2_dev\externals\jitify/jitify.hpp(2866): here
  
  3 errors detected in the compilation of "C:/Users/appveyor/AppData/Local/Temp/1/tmpxft_00000d18_00000000-8_CUDAScanCompaction.compute_60.cpp1.ii".
  CUDAScanCompaction.cu

https://ci.appveyor.com/project/mondus/flamegpu2-dev/builds/32555760#L182