[cfe-commits] r121653 - in /cfe/trunk: include/clang/Basic/DiagnosticSemaKinds.td lib/Sema/SemaDeclAttr.cpp test/SemaCUDA/qualifiers.cu

[Peter Collingbourne]

Author: pcc
Date: Sun Dec 12 17:02:57 2010
New Revision: 121653

URL: http://llvm.org/viewvc/llvm-project?rev=121653&view=rev
Log:
Sema: diagnose kernel functions with non-void return type

Modified:
    cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
    cfe/trunk/lib/Sema/SemaDeclAttr.cpp
    cfe/trunk/test/SemaCUDA/qualifiers.cu

Modified: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td?rev=121653&r1=121652&r2=121653&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td (original)
+++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td Sun Dec 12 17:02:57 2010
@@ -2859,9 +2859,12 @@
   "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">;

Modified: cfe/trunk/lib/Sema/SemaDeclAttr.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDeclAttr.cpp?rev=121653&r1=121652&r2=121653&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaDeclAttr.cpp (original)
+++ cfe/trunk/lib/Sema/SemaDeclAttr.cpp Sun Dec 12 17:02:57 2010
@@ -2168,6 +2168,21 @@
       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";

Modified: cfe/trunk/test/SemaCUDA/qualifiers.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/qualifiers.cu?rev=121653&r1=121652&r2=121653&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/qualifiers.cu (original)
+++ cfe/trunk/test/SemaCUDA/qualifiers.cu Sun Dec 12 17:02:57 2010
@@ -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;
+}