Sema: diagnose kernel functions with non-void return type

llvm-svn: 121653

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -2859,9 +2859,12 @@ def err_atomic_builtin_pointer_size : Error<
   "first argument to atomic builtin must be a pointer to 1,2,4,8 or 16 byte "
   "type (%0 invalid)">;
 
-  
 def err_deleted_function_use : Error<"attempt to use a deleted function">;
 
+def err_kern_type_not_void_return : Error<
+  "kernel function type %0 must have void return type">;
+
+
 def err_cannot_pass_objc_interface_to_vararg : Error<
   "cannot pass object with interface type %0 by-value through variadic "
   "%select{function|block|method}1">;

clang/lib/Sema/SemaDeclAttr.cpp

@@ -2168,6 +2168,21 @@ static void HandleGlobalAttr(Decl *d, const AttributeList &Attr, Sema &S) {
       return;
     }
 
+    FunctionDecl *FD = cast<FunctionDecl>(d);
+    if (!FD->getResultType()->isVoidType()) {
+      TypeLoc TL = FD->getTypeSourceInfo()->getTypeLoc();
+      if (FunctionTypeLoc* FTL = dyn_cast<FunctionTypeLoc>(&TL)) {
+        S.Diag(FD->getTypeSpecStartLoc(), diag::err_kern_type_not_void_return)
+          << FD->getType()
+          << FixItHint::CreateReplacement(FTL->getResultLoc().getSourceRange(),
+                                          "void");
+      } else {
+        S.Diag(FD->getTypeSpecStartLoc(), diag::err_kern_type_not_void_return)
+          << FD->getType();
+      }
+      return;
+    }
+
     d->addAttr(::new (S.Context) CUDAGlobalAttr(Attr.getLoc(), S.Context));
   } else {
     S.Diag(Attr.getLoc(), diag::warn_attribute_ignored) << "global";

clang/test/SemaCUDA/qualifiers.cu

@@ -3,3 +3,6 @@
 #include "cuda.h"
 
 __global__ void g1(int x) {}
+__global__ int g2(int x) { // expected-error {{must have void return type}}
+  return 1;
+}