[cfe-dev] Clang and CUDA with C++11 features

[Peter Colberg]

On Mon, Jun 18, 2012 at 04:44:05PM -0700, Richard Smith wrote:
> On Mon, Jun 18, 2012 at 3:46 PM, Peter Colberg <peter at colberg.org> wrote:
> > When including <iostream>, there is still a minor problem:
> >
> >    // template.cu
> >
> >    #include <cuda_runtime.h>
> >    #include <iostream>
> >
> >    int main() {}
> >
> >
> >    clang++ -Wall -I/usr/local/cuda-4.2/cuda/include -L/usr/local/cuda-4.2/cuda/lib64 -lcudart -o template template.cu
> [...]
> >    In file included from template.cu:2:
> >    In file included from /home/peter/usr/rhel6-x86_64/gcc-4.7.1/lib/gcc/x86_64-unknown-linux-gnu/4.7.1/../../../../include/c++/4.7.1/iostream:39:
> >    In file included from /home/peter/usr/rhel6-x86_64/gcc-4.7.1/lib/gcc/x86_64-unknown-linux-gnu/4.7.1/../../../../include/c++/4.7.1/ostream:39:
> >    In file included from /home/peter/usr/rhel6-x86_64/gcc-4.7.1/lib/gcc/x86_64-unknown-linux-gnu/4.7.1/../../../../include/c++/4.7.1/ios:42:
> >    In file included from /home/peter/usr/rhel6-x86_64/gcc-4.7.1/lib/gcc/x86_64-unknown-linux-gnu/4.7.1/../../../../include/c++/4.7.1/bits/ios_base.h:42:
> >    In file included from /home/peter/usr/rhel6-x86_64/gcc-4.7.1/lib/gcc/x86_64-unknown-linux-gnu/4.7.1/../../../../include/c++/4.7.1/bits/locale_classes.h:41:
> >    In file included from /home/peter/usr/rhel6-x86_64/gcc-4.7.1/lib/gcc/x86_64-unknown-linux-gnu/4.7.1/../../../../include/c++/4.7.1/string:53:
> >    /home/peter/usr/rhel6-x86_64/gcc-4.7.1/lib/gcc/x86_64-unknown-linux-gnu/4.7.1/../../../../include/c++/4.7.1/bits/basic_string.h:504:9: error: call to
> >          implicitly-deleted copy constructor of 'std::basic_string<char, std::char_traits<char>, std::allocator<char>>::_Alloc_hider'
> >          : _M_dataplus(__str._M_dataplus)
> >            ^           ~~~~~~~~~~~~~~~~~
> >    /home/peter/usr/rhel6-x86_64/gcc-4.7.1/lib/gcc/x86_64-unknown-linux-gnu/4.7.1/../../../../include/c++/4.7.1/bits/basic_string.h:2863:12: note: in
> >          instantiation of member function 'std::basic_string<char, std::char_traits<char>, std::allocator<char> >::basic_string' requested here
> >      { return __gnu_cxx::__to_xstring<string>(&std::vsnprintf, 4 * sizeof(int),
> >               ^
> >    /home/peter/usr/rhel6-x86_64/gcc-4.7.1/lib/gcc/x86_64-unknown-linux-gnu/4.7.1/../../../../include/c++/4.7.1/bits/basic_string.h:268:29: note: copy
> >          constructor of '_Alloc_hider' is implicitly deleted because base class 'std::allocator<char>' has no copy constructor
> >          struct _Alloc_hider : _Alloc
> >                                ^
> 
> This "has no copy constructor" condition should not be possible. Looks
> like the problem is here (end of Sema/SemaDeclCXX.cpp):
> 
> Sema::CUDAFunctionTarget Sema::IdentifyCUDATarget(const FunctionDecl *D) {
>   // Implicitly declared functions (e.g. copy constructors) are
>   // __host__ __device__
>   if (D->isImplicit())
>     return CFT_HostDevice;
> 
> bool Sema::CheckCUDATarget(CUDAFunctionTarget CallerTarget,
>                            CUDAFunctionTarget CalleeTarget) {
>   if (CallerTarget == CFT_HostDevice && CalleeTarget != CFT_HostDevice)
>     return true;
> 
> So... an implicit copy constructor is __host__ __device__, and thus
> can't call an explicitly-declared copy constructor (which is just
> __host__). Oops. Is this a bug in the CUDA spec or in Clang?

Please correct me if I am wrong, but I have never seen a CUDA
specification. The CUDA programming guide does not mention
copy constructors.

Why are these attributes needed at all? Can't the compiler just fail
on code that will not compile to a GPU kernel, e.g. due to use of C++
run-time support?

I have never understood why the nvcc compiler rejects perfectly
valid C++ code simply because of a nonexistent __device__ attribute.

Peter