[cfe-commits] r126292 - in /cfe/trunk: include/clang/Basic/DiagnosticSemaKinds.td lib/Sema/SemaExpr.cpp test/SemaCUDA/kernel-call.cu

[Peter Collingbourne]

Author: pcc
Date: Tue Feb 22 19:53:29 2011
New Revision: 126292

URL: http://llvm.org/viewvc/llvm-project?rev=126292&view=rev
Log:
Sema: diagnose kernel calls to non-global functions

Modified:
    cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
    cfe/trunk/lib/Sema/SemaExpr.cpp
    cfe/trunk/test/SemaCUDA/kernel-call.cu

Modified: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td?rev=126292&r1=126291&r2=126292&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td (original)
+++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td Tue Feb 22 19:53:29 2011
@@ -3121,6 +3121,8 @@
   "kernel function type %0 must have void return type">;
 def err_config_scalar_return : Error<
   "CUDA special function 'cudaConfigureCall' must have scalar return type">;
+def err_kern_call_not_global_function : Error<
+  "kernel call to non-global function %0">;
 
 
 def err_cannot_pass_objc_interface_to_vararg : Error<

Modified: cfe/trunk/lib/Sema/SemaExpr.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaExpr.cpp?rev=126292&r1=126291&r2=126292&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaExpr.cpp (original)
+++ cfe/trunk/lib/Sema/SemaExpr.cpp Tue Feb 22 19:53:29 2011
@@ -4625,6 +4625,20 @@
     return ExprError(Diag(LParenLoc, diag::err_typecheck_call_not_function)
       << Fn->getType() << Fn->getSourceRange());
 
+  if (getLangOptions().CUDA) {
+    if (Config) {
+      // CUDA: Kernel calls must be to global functions
+      if (FDecl && !FDecl->hasAttr<CUDAGlobalAttr>())
+        return ExprError(Diag(LParenLoc,diag::err_kern_call_not_global_function)
+            << FDecl->getName() << Fn->getSourceRange());
+
+      // CUDA: Kernel function must have 'void' return type
+      if (!FuncT->getResultType()->isVoidType())
+        return ExprError(Diag(LParenLoc, diag::err_kern_type_not_void_return)
+            << Fn->getType() << Fn->getSourceRange());
+    }
+  }
+
   // Check for a valid return type
   if (CheckCallReturnType(FuncT->getResultType(),
                           Fn->getSourceRange().getBegin(), TheCall,

Modified: cfe/trunk/test/SemaCUDA/kernel-call.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/kernel-call.cu?rev=126292&r1=126291&r2=126292&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/kernel-call.cu (original)
+++ cfe/trunk/test/SemaCUDA/kernel-call.cu Tue Feb 22 19:53:29 2011
@@ -8,8 +8,16 @@
   g1<<<arg, arg>>>(1);
 }
 
+void h1(int x) {}
+int h2(int x) { return 1; }
+
 int main(void) {
   g1<<<1, 1>>>(42);
 
   t1(1);
+
+  h1<<<1, 1>>>(42); // expected-error {{kernel call to non-global function h1}}
+
+  int (*fp)(int) = h2;
+  fp<<<1, 1>>>(42); // expected-error {{must have void return type}}
 }