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

Peter Colberg peter at colberg.org
Mon Jun 18 15:46:38 PDT 2012


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:1:
    /usr/local/cuda-4.2/cuda/include/cuda_runtime.h:267:10: warning: function 'cudaMallocHost' has internal linkage but is used in an inline function with external
          linkage [-Winternal-linkage-in-inline]
      return cudaMallocHost((void**)(void*)ptr, size, flags);
             ^
    /usr/local/cuda-4.2/cuda/include/cuda_runtime.h:261:1: note: use 'static' to give inline function 'cudaMallocHost' internal linkage
    __inline__ __host__ cudaError_t cudaMallocHost(
    ^
    static 
    /usr/local/cuda-4.2/cuda/include/cuda_runtime.h:222:40: note: 'cudaMallocHost' declared here
    static __inline__ __host__ cudaError_t cudaMallocHost(
                                           ^
    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
                                ^
    /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<wchar_t, std::char_traits<wchar_t>, std::allocator<wchar_t>>::_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:2964:12: note: in
          instantiation of member function 'std::basic_string<wchar_t, std::char_traits<wchar_t>, std::allocator<wchar_t> >::basic_string' requested here
      { return __gnu_cxx::__to_xstring<wstring>(&std::vswprintf, 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<wchar_t>' has no copy constructor
          struct _Alloc_hider : _Alloc
                                ^
    1 warning and 2 errors generated.


Is there something else that needs to be added for C++11 parsing,
besides CPlusPlus0x in include/clang/Frontend/LangStandards.def?

The same compiles fine in C++ mode (modulo warning):

    // template.cpp

    #include <cuda_runtime.h>
    #include <iostream>

    int main() {}


    clang++ -std=c++11 -Wall -I/usr/local/cuda-4.2/cuda/include -L/usr/local/cuda-4.2/cuda/lib64 -lcudart -o template template.cpp
    In file included from template.cpp:1:
    /usr/local/cuda-4.2/cuda/include/cuda_runtime.h:267:10: warning: function 'cudaMallocHost' has internal linkage but is used in an inline function with external
          linkage [-Winternal-linkage-in-inline]
      return cudaMallocHost((void**)(void*)ptr, size, flags);
             ^
    /usr/local/cuda-4.2/cuda/include/cuda_runtime.h:261:1: note: use 'static' to give inline function 'cudaMallocHost' internal linkage
    __inline__ __host__ cudaError_t cudaMallocHost(
    ^
    static 
    /usr/local/cuda-4.2/cuda/include/cuda_runtime.h:222:40: note: 'cudaMallocHost' declared here
    static __inline__ __host__ cudaError_t cudaMallocHost(
                                           ^
    1 warning generated.

Regards,
Peter



More information about the cfe-dev mailing list