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

[Richard Smith]

On Mon, Jun 18, 2012 at 3:46 PM, Peter Colberg <peter at colberg.org> wrote:
> On Sun, Jun 17, 2012 at 11:22:58PM -0700, Richard Smith wrote:
>> I've just improved the diagnostic for this in Clang ToT:
>>
>> test/Parser/cuda-kernel-call.cu:13:12: error: a space is required between
>>       consecutive right angle brackets (use '> >')
>>   S<S<S<int>>> s; // expected-error 2{{use '> >'}}
>>            ^~
>>            > >
>> test/Parser/cuda-kernel-call.cu:13:13: error: a space is required between
>>       consecutive right angle brackets (use '> >')
>>   S<S<S<int>>> s; // expected-error 2{{use '> >'}}
>>             ^~
>>             > >
>>
>> If we want to support CUDA + C++11, I think the right fix for this
>> issue is to extend C++11's rule for '>>'-fission to CUDA's '>>>'
>> token:
>>
>> Index: lib/Parse/ParseTemplate.cpp
>> ===================================================================
>> --- lib/Parse/ParseTemplate.cpp (revision 158652)
>> +++ lib/Parse/ParseTemplate.cpp (working copy)
>> @@ -785,7 +785,8 @@
>>      Hint2 = FixItHint::CreateInsertion(Next.getLocation(), " ");
>>
>>    unsigned DiagId = diag::err_two_right_angle_brackets_need_space;
>> -  if (getLangOpts().CPlusPlus0x && Tok.is(tok::greatergreater))
>> +  if (getLangOpts().CPlusPlus0x &&
>> +      (Tok.is(tok::greatergreater) || Tok.is(tok::greatergreatergreater)))
>>      DiagId = diag::warn_cxx98_compat_two_right_angle_brackets;
>>    else if (Tok.is(tok::greaterequal))
>>      DiagId = diag::err_right_angle_bracket_equal_needs_space;
>>
>> This will allow us to correctly deal with '>>' in almost all cases in
>> CUDA mode, and the exceptions are truly bizarre and implausible
>> constructs like "(SomeType)&FnTmpl<ClassTmpl<T>>>>3;" -- here the
>> '>>>>' splits into '> > >>' in C++11 and would split into '> > > >' in
>> CUDA/C++11 with the proposed rule.
>
> With the above patch, C++11 template parameters and CUDA syntax coexist nicely:
>
>    #include <cuda_runtime.h>
>
>    template <typename T>
>    struct S;
>
>    template <typename T>
>    __attribute__((global)) void f() {}
>
>    int main()
>    {
>        f<S<S<int>>><<<1, 1>>>();
>    }
>
> 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?