Author: jlebar Date: Tue Jan 19 18:26:57 2016 New Revision: 258263 URL: http://llvm.org/viewvc/llvm-project?rev=258263&view=rev Log: [CUDA] Only allow __global__ on free functions and static member functions.
Summary: Warn for NVCC compatibility if you declare a static member function or inline function as __global__. Reviewers: tra Subscribers: jhen, echristo, cfe-commits Differential Revision: http://reviews.llvm.org/D16261 Modified: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td cfe/trunk/lib/Sema/SemaCUDA.cpp cfe/trunk/lib/Sema/SemaDeclAttr.cpp cfe/trunk/test/SemaCUDA/bad-attributes.cu Modified: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td?rev=258263&r1=258262&r2=258263&view=diff ============================================================================== --- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td (original) +++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td Tue Jan 19 18:26:57 2016 @@ -6403,6 +6403,8 @@ def err_deleted_function_use : Error<"at def err_kern_type_not_void_return : Error< "kernel function type %0 must have void return type">; +def err_kern_is_nonstatic_method : Error< + "kernel function %0 must be a free function or static member function">; def err_config_scalar_return : Error< "CUDA special function 'cudaConfigureCall' must have scalar return type">; def err_kern_call_not_global_function : Error< @@ -6415,6 +6417,12 @@ def err_ref_bad_target : Error< def warn_host_calls_from_host_device : Warning< "calling __host__ function %0 from __host__ __device__ function %1 can lead to runtime errors">, InGroup<CudaCompat>; +def warn_kern_is_method : Extension< + "kernel function %0 is a member function; this may not be accepted by nvcc">, + InGroup<CudaCompat>; +def warn_kern_is_inline : Warning< + "ignored 'inline' attribute on kernel function %0">, + InGroup<CudaCompat>; def warn_non_pod_vararg_with_format_string : Warning< "cannot pass %select{non-POD|non-trivial}0 object of type %1 to variadic " Modified: cfe/trunk/lib/Sema/SemaCUDA.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaCUDA.cpp?rev=258263&r1=258262&r2=258263&view=diff ============================================================================== --- cfe/trunk/lib/Sema/SemaCUDA.cpp (original) +++ cfe/trunk/lib/Sema/SemaCUDA.cpp Tue Jan 19 18:26:57 2016 @@ -273,12 +273,9 @@ static bool resolveCalleeCUDATargetConflict(Sema::CUDAFunctionTarget Target1, Sema::CUDAFunctionTarget Target2, Sema::CUDAFunctionTarget *ResolvedTarget) { - if (Target1 == Sema::CFT_Global && Target2 == Sema::CFT_Global) { - // TODO: this shouldn't happen, really. Methods cannot be marked __global__. - // Clang should detect this earlier and produce an error. Then this - // condition can be changed to an assertion. - return true; - } + // Only free functions and static member functions may be global. + assert(Target1 != Sema::CFT_Global); + assert(Target2 != Sema::CFT_Global); if (Target1 == Sema::CFT_HostDevice) { *ResolvedTarget = Target2; Modified: cfe/trunk/lib/Sema/SemaDeclAttr.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDeclAttr.cpp?rev=258263&r1=258262&r2=258263&view=diff ============================================================================== --- cfe/trunk/lib/Sema/SemaDeclAttr.cpp (original) +++ cfe/trunk/lib/Sema/SemaDeclAttr.cpp Tue Jan 19 18:26:57 2016 @@ -3617,11 +3617,21 @@ static void handleGlobalAttr(Sema &S, De : FixItHint()); return; } + if (const auto *Method = dyn_cast<CXXMethodDecl>(FD)) { + if (Method->isInstance()) { + S.Diag(Method->getLocStart(), diag::err_kern_is_nonstatic_method) + << Method; + return; + } + S.Diag(Method->getLocStart(), diag::warn_kern_is_method) << Method; + } + // Only warn for "inline" when compiling for host, to cut down on noise. + if (FD->isInlineSpecified() && !S.getLangOpts().CUDAIsDevice) + S.Diag(FD->getLocStart(), diag::warn_kern_is_inline) << FD; D->addAttr(::new (S.Context) CUDAGlobalAttr(Attr.getRange(), S.Context, Attr.getAttributeSpellingListIndex())); - } static void handleGNUInlineAttr(Sema &S, Decl *D, const AttributeList &Attr) { Modified: cfe/trunk/test/SemaCUDA/bad-attributes.cu URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/bad-attributes.cu?rev=258263&r1=258262&r2=258263&view=diff ============================================================================== --- cfe/trunk/test/SemaCUDA/bad-attributes.cu (original) +++ cfe/trunk/test/SemaCUDA/bad-attributes.cu Tue Jan 19 18:26:57 2016 @@ -4,8 +4,8 @@ // // You should be able to run this file through nvcc for compatibility testing. // -// RUN: %clang_cc1 -fsyntax-only -verify %s -// RUN: %clang_cc1 -fcuda-is-device -fsyntax-only -verify %s +// RUN: %clang_cc1 -fsyntax-only -Wcuda-compat -verify -DEXPECT_INLINE_WARNING %s +// RUN: %clang_cc1 -fcuda-is-device -fsyntax-only -Wcuda-compat -verify %s #include "Inputs/cuda.h" @@ -47,3 +47,15 @@ __global__ __device__ void z11(); // ex // expected-note@-1 {{conflicting attribute is here}} __global__ __host__ void z12(); // expected-error {{attributes are not compatible}} // expected-note@-1 {{conflicting attribute is here}} + +struct S { + __global__ void foo() {}; // expected-error {{must be a free function or static member function}} + __global__ static void bar(); // expected-warning {{kernel function 'bar' is a member function}} + // Although this is implicitly inline, we shouldn't warn. + __global__ static void baz() {}; // expected-warning {{kernel function 'baz' is a member function}} +}; + +__global__ static inline void foobar() {}; +#ifdef EXPECT_INLINE_WARNING +// expected-warning@-2 {{ignored 'inline' attribute on kernel function 'foobar'}} +#endif _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
