This revision was automatically updated to reflect the committed changes.
jlebar marked 3 inline comments as done.
Closed by commit rL258263: [CUDA] Only allow __global__ on free functions and
static member functions. (authored by jlebar).
Changed prior to commit:
http://reviews.llvm.org/D16261?vs=45297&id=45324#toc
Repository:
rL LLVM
http://reviews.llvm.org/D16261
Files:
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
Index: cfe/trunk/lib/Sema/SemaCUDA.cpp
===================================================================
--- cfe/trunk/lib/Sema/SemaCUDA.cpp
+++ cfe/trunk/lib/Sema/SemaCUDA.cpp
@@ -273,12 +273,9 @@
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;
Index: cfe/trunk/lib/Sema/SemaDeclAttr.cpp
===================================================================
--- cfe/trunk/lib/Sema/SemaDeclAttr.cpp
+++ cfe/trunk/lib/Sema/SemaDeclAttr.cpp
@@ -3617,11 +3617,21 @@
: 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) {
Index: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
+++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
@@ -6403,6 +6403,8 @@
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 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 "
Index: cfe/trunk/test/SemaCUDA/bad-attributes.cu
===================================================================
--- cfe/trunk/test/SemaCUDA/bad-attributes.cu
+++ cfe/trunk/test/SemaCUDA/bad-attributes.cu
@@ -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 @@
// 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
