[PATCH] D16261: [CUDA] Only allow __global__ on free functions and static member functions.

[Justin Lebar via cfe-commits]

jlebar created this revision.
jlebar added a reviewer: tra.
jlebar added subscribers: cfe-commits, echristo, jhen.

Warn for NVCC compatibility if you declare a static member function or
inline function as __global__.

http://reviews.llvm.org/D16261

Files:
  include/clang/Basic/DiagnosticSemaKinds.td
  lib/Sema/SemaCUDA.cpp
  lib/Sema/SemaDeclAttr.cpp
  test/SemaCUDA/bad-attributes.cu

Index: test/SemaCUDA/bad-attributes.cu
===================================================================
--- test/SemaCUDA/bad-attributes.cu
+++ test/SemaCUDA/bad-attributes.cu
@@ -47,3 +47,12 @@
 // expected-note at -1 {{conflicting attribute is here}}
 __global__ __host__ void z12();  // expected-error {{attributes are not compatible}}
 // expected-note at -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}}
+  __global__ static void baz() {}; // expected-warning {{kernel function 'baz' is a member function}}
+  // expected-warning at -1 {{kernel function 'baz' is inline}}
+};
+
+__global__ static inline void foobar() {};  // expected-warning {{kernel function 'foobar' is inline}}
Index: lib/Sema/SemaDeclAttr.cpp
===================================================================
--- lib/Sema/SemaDeclAttr.cpp
+++ lib/Sema/SemaDeclAttr.cpp
@@ -3617,11 +3617,23 @@
                               : 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_nvcc_compat_kern_is_method)
+        << Method;
+  }
+  if (FD->isInlined()) {
+    S.Diag(FD->getLocStart(), diag::warn_nvcc_compat_kern_is_inlined)
+        << FD;
+  }
 
   D->addAttr(::new (S.Context)
               CUDAGlobalAttr(Attr.getRange(), S.Context,
                              Attr.getAttributeSpellingListIndex()));
-
 }
 
 static void handleGNUInlineAttr(Sema &S, Decl *D, const AttributeList &Attr) {
Index: lib/Sema/SemaCUDA.cpp
===================================================================
--- lib/Sema/SemaCUDA.cpp
+++ 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: include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- include/clang/Basic/DiagnosticSemaKinds.td
+++ include/clang/Basic/DiagnosticSemaKinds.td
@@ -6400,6 +6400,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<
@@ -6412,6 +6414,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_nvcc_compat_kern_is_method : Warning<
+  "kernel function %0 is a member function; this may not be accepted by nvcc">,
+  InGroup<CudaCompat>;
+def warn_nvcc_compat_kern_is_inlined : Warning<
+  "kernel function %0 is inlined; this may not be accepted by nvcc">,
+  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 "


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D16261.45080.patch
Type: text/x-patch
Size: 3963 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20160116/b2ad97bf/attachment.bin>