[NV 200636681] Compile Time Failure "implicit conversion from 'const uint_type' to 'float'" with Clang 10.0.0/gpgpu/cuda_a on Ubuntu18.04/x86_64

[alliepiper]

#$ cicc --c++14 --clang --clang_version=100000 --promote_warnings --orig_src_file_name "/dvs/p4/build/sw/gpgpu/thrust/testing//set_intersection.cu" --unicode_source_kind=UTF-8 --allow_managed  --display_error_number  -arch compute_35 -m64 --no-version-ident -ftz=0 -prec_div=1 -prec_sqrt=1 -fmad=1 --include_file_name "set_intersection.fatbin.c" -tused -nvvmir-library "/dvs/p4/build/sw/gpgpu/bin/x86_64_Linux_release/nvvm/libdevice/libdevice.10.bc" --module_id_file_name "set_intersection.o.keep_dir/set_intersection.module_id" --gen_c_file_name "set_intersection.o.keep_dir/set_intersection.compute_35.cudafe1.c" --stub_file_name "set_intersection.o.keep_dir/set_intersection.compute_35.cudafe1.stub.c" --gen_device_file_name "set_intersection.o.keep_dir/set_intersection.compute_35.cudafe1.gpu"  "set_intersection.o.keep_dir/set_intersection.compute_35.cpp1.ii" -o "set_intersection.o.keep_dir/set_intersection.compute_35.ptx"
/dvs/p4/build/sw/gpgpu/thrust/thrust/random/detail/normal_distribution_base.h:53:51: error: implicit conversion from 'const uint_type' (aka 'const unsigned int') to 'float' changes value from 2147483645 to 2147483648 [-Werror,-Wimplicit-int-float-conversion]
const RealType S1 = (static_cast< RealType>(1)) / urng_range; 
                                                ~ ^~~~~~~~~~
/dvs/p4/build/sw/gpgpu/thrust/thrust/random/detail/normal_distribution.inl:83:17: note: in instantiation of function template specialization 'thrust::random::detail::normal_distribution_nvcc<float>::sample<thrust::random::linear_congruential_engine<unsigned int, 48271, 0, 2147483647> >' requested here
return super_t::sample(urng, (parm.first), (parm.second)); 

NVBug 200636681

[alliepiper]

Starting working through clang 10 issues here: https://github.com/allisonvacanti/thrust/tree/bug/clang10_fixes/gh.1268/nv.200636681

Currently running into what appears to be a compiler bug in the unordered_map implementation used by per_device_resource that will need to be worked around.

/usr/lib/gcc/x86_64-linux-gnu/10/../../../../include/c++/10/tuple(566): error #1921: pack "_UElements" does not have the same number of elements as "_Elements"
          detected during:
            instantiation of "__nv_bool std::tuple<_Elements...>::__nothrow_constructible<_UElements...>() [with _Elements=<const int &>, _UElements=<>]"
/usr/lib/gcc/x86_64-linux-gnu/10/../../../../include/c++/10/bits/hashtable_policy.h(715): here
            instantiation of "std::__detail::_Map_base<_Key, _Pair, _Alloc, std::__detail::_Select1st, _Equal, _H1, _H2, _Hash, _RehashPolicy, _Traits, true>::mapped_type &std::__detail::_Map_base<_Key, _Pair, _Alloc, std::__detail::_Select1st, _Equal, _H1, _H2, _Hash, _RehashPolicy, _Traits, true>::operator[](const std::__detail::_Map_base<_Key, _Pair, _Alloc, std::__detail::_Select1st, _Equal, _H1, _H2, _Hash, _RehashPolicy, _Traits, true>::key_type &) [with _Key=int, _Pair=std::pair<const int, thrust::system::cuda::detail::default_async_device_resource>, _Alloc=std::allocator<std::pair<const int, thrust::system::cuda::detail::default_async_device_resource>>, _Equal=std::equal_to<int>, _H1=std::hash<int>, _H2=std::__detail::_Mod_range_hashing, _Hash=std::__detail::_Default_ranged_hash, _RehashPolicy=std::__detail::_Prime_rehash_policy, _Traits=std::__umap_traits<false>]"
/usr/lib/gcc/x86_64-linux-gnu/10/../../../../include/c++/10/bits/unordered_map.h(984): here
            instantiation of "std::unordered_map<_Key, _Tp, _Hash, _Pred, _Alloc>::mapped_type &std::unordered_map<_Key, _Tp, _Hash, _Pred, _Alloc>::operator[](const std::unordered_map<_Key, _Tp, _Hash, _Pred, _Alloc>::key_type &) [with _Key=int, _Tp=thrust::system::cuda::detail::default_async_device_resource, _Hash=std::hash<int>, _Pred=std::equal_to<int>, _Alloc=std::allocator<std::pair<const int, thrust::system::cuda::detail::default_async_device_resource>>]"
/home/av/code/src/thrust/thrust/system/cuda/detail/per_device_resource.h(63): here
            instantiation of "MR *thrust::cuda_cub::get_per_device_resource<MR,DerivedPolicy>(thrust::cuda_cub::execution_policy<DerivedPolicy> &) [with MR=thrust::system::cuda::detail::default_async_device_resource, DerivedPolicy=thrust::cuda_cub::par_t]"
/home/av/code/src/thrust/thrust/per_device_resource.h(48): here
            instantiation of "MR *thrust::get_per_device_resource<MR,DerivedPolicy>(const thrust::detail::execution_policy_base<DerivedPolicy> &) [with MR=thrust::system::cuda::detail::default_async_device_resource, DerivedPolicy=thrust::cuda_cub::par_t]"
/home/av/code/src/thrust/thrust/per_device_resource.h(81): here
            instantiation of "thrust::per_device_allocator<T, Upstream, ExecutionPolicy>::per_device_allocator() [with T=thrust::detail::uint8_t, Upstream=thrust::system::cuda::detail::default_async_device_resource, ExecutionPolicy=thrust::cuda_cub::par_t]"
/home/av/code/src/thrust/testing/async_reduce_into.cu(401): here
            instantiation of "void test_async_reduce_into<AsyncReduceIntoInvoker, SyncReduceIntoInvoker>::tester<T>::operator()(std::size_t) [with AsyncReduceIntoInvoker=reduce_into_async_invoker, SyncReduceIntoInvoker=reduce_sync_invoker, T=char]

[alliepiper]

That error seems to happen for many usages of the STL maps with nvcc 11.1 / Clang 10 / libc++ 10 / c++17. It does not happen for c++14 or c++11.

$ cat t.cu
#include <unordered_map>

template<typename MR>
__host__
MR * get_per_device_resource(int k)
{
  static std::unordered_map<int, MR> device_id_to_resource;
  return &device_id_to_resource[k];
}

float* bar(int k)
{
  return get_per_device_resource<float>(k);
}

av@torsey /tmp$ nvcc -ccbin=/usr/bin/clang++-10 --std=c++17 -c t.cu
/usr/bin/../lib/gcc/x86_64-linux-gnu/10/../../../../include/c++/10/tuple(566): error: pack "_UElements" does not have the same number of elements as "_Elements"
          detected during:
            instantiation of "__nv_bool std::tuple<_Elements...>::__nothrow_constructible<_UElements...>() [with _Elements=<const int &>, _UElements=<>]"
/usr/bin/../lib/gcc/x86_64-linux-gnu/10/../../../../include/c++/10/bits/hashtable_policy.h(715): here
            instantiation of "std::__detail::_Map_base<_Key, _Pair, _Alloc, std::__detail::_Select1st, _Equal, _H1, _H2, _Hash, _RehashPolicy, _Traits, true>::mapped_type &std::__detail::_Map_base<_Key, _Pair, _Alloc, std::__detail::_Select1st, _Equal, _H1, _H2, _Hash, _RehashPolicy, _Traits, true>::operator[](const std::__detail::_Map_base<_Key, _Pair, _Alloc, std::__detail::_Select1st, _Equal, _H1, _H2, _Hash, _RehashPolicy, _Traits, true>::key_type &) [with _Key=int, _Pair=std::pair<const int, float>, _Alloc=std::allocator<std::pair<const int, float>>, _Equal=std::equal_to<int>, _H1=std::hash<int>, _H2=std::__detail::_Mod_range_hashing, _Hash=std::__detail::_Default_ranged_hash, _RehashPolicy=std::__detail::_Prime_rehash_policy, _Traits=std::__umap_traits<false>]"
/usr/bin/../lib/gcc/x86_64-linux-gnu/10/../../../../include/c++/10/bits/unordered_map.h(984): here
            instantiation of "std::unordered_map<_Key, _Tp, _Hash, _Pred, _Alloc>::mapped_type &std::unordered_map<_Key, _Tp, _Hash, _Pred, _Alloc>::operator[](const std::unordered_map<_Key, _Tp, _Hash, _Pred, _Alloc>::key_type &) [with _Key=int, _Tp=float, _Hash=std::hash<int>, _Pred=std::equal_to<int>, _Alloc=std::allocator<std::pair<const int, float>>]"
t.cu(8): here
            instantiation of "MR *get_per_device_resource<MR>(int) [with MR=float]"
t.cu(13): here

1 error detected in the compilation of "t.cu".

$ nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2020 NVIDIA Corporation
Built on Tue_Sep_15_19:10:02_PDT_2020
Cuda compilation tools, release 11.1, V11.1.74
Build cuda_11.1.TC455_06.29069683_0

$ /usr/bin/clang++-10 --version
clang version 10.0.0-4ubuntu1~18.04.2
Target: x86_64-pc-linux-gnu
Thread model: posix
InstalledDir: /usr/bin

# Works without nvcc:
$ /usr/bin/clang++-10 -D__host__= --std=c++17 -x c++ -c t.cu; echo $?
0

Filed NVBug 3135576, this issue is blocked pending fixes/workarounds from the compiler folks.