Sema: diagnose kernel functions with non-void return type

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@121653 91177308-0d34-0410-b5e6-96231b3b80d8

diff --git a/include/clang/Basic/DiagnosticSemaKinds.td b/include/clang/Basic/DiagnosticSemaKinds.td
index 55270b0116f52d36d475794dd56cdbcd3e33be81..1b23aaaf638926964f61996a3f7d51c951b342be 100644 (file)
--- a/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/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">;

diff --git a/lib/Sema/SemaDeclAttr.cpp b/lib/Sema/SemaDeclAttr.cpp
index e5f27fcfbb88e2458dacbc152c0c6ecc243fd84f..a2a7fdcdd5392505704c5227ad34a39e1676bb49 100644 (file)
--- a/lib/Sema/SemaDeclAttr.cpp
+++ b/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";

diff --git a/test/SemaCUDA/qualifiers.cu b/test/SemaCUDA/qualifiers.cu
index 8d5b759a6d7bc4879007ab3ee269abb92dc4ea26..1346d654b8c419f54f02c089af56be571d2ffa46 100644 (file)
--- a/test/SemaCUDA/qualifiers.cu
+++ b/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;
+}