commit 8591a7f5215f6469603e6bc8f4fdbcc78e215ab9
author Peter Collingbourne <peter@pcc.me.uk> Sun Oct 02 23:49:15 2011 +0000
committer Peter Collingbourne <peter@pcc.me.uk> Sun Oct 02 23:49:15 2011 +0000
tree e99806082ae238b244b91942709764bc87ab3928
parent 80d2355df4a87b52412cc7481496c4a4eccd3db1

CUDA: diagnose unconfigured calls to global functions

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

include/clang/Basic/DiagnosticSemaKinds.td

diff --git a/include/clang/Basic/DiagnosticSemaKinds.td b/include/clang/Basic/DiagnosticSemaKinds.td
index 14b36e4..357f1af 100644
--- a/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/include/clang/Basic/DiagnosticSemaKinds.td
@@ -3986,6 +3986,8 @@
   "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_global_call_not_config : Error<
+  "call to global function %0 not configured">;
 
 
 def err_cannot_pass_objc_interface_to_vararg : Error<

lib/Sema/SemaExpr.cpp

diff --git a/lib/Sema/SemaExpr.cpp b/lib/Sema/SemaExpr.cpp
index d0a3f6b..eb5678e 100644
--- a/lib/Sema/SemaExpr.cpp
+++ b/lib/Sema/SemaExpr.cpp
@@ -3655,6 +3655,11 @@
       if (!FuncT->getResultType()->isVoidType())
         return ExprError(Diag(LParenLoc, diag::err_kern_type_not_void_return)
             << Fn->getType() << Fn->getSourceRange());
+    } else {
+      // CUDA: Calls to global functions must be configured
+      if (FDecl && FDecl->hasAttr<CUDAGlobalAttr>())
+        return ExprError(Diag(LParenLoc, diag::err_global_call_not_config)
+            << FDecl->getName() << Fn->getSourceRange());
     }
   }
 

test/SemaCUDA/kernel-call.cu

diff --git a/test/SemaCUDA/kernel-call.cu b/test/SemaCUDA/kernel-call.cu
index 7bc7ae1..9a22aa7 100644
--- a/test/SemaCUDA/kernel-call.cu
+++ b/test/SemaCUDA/kernel-call.cu
@@ -13,6 +13,7 @@
 
 int main(void) {
   g1<<<1, 1>>>(42);
+  g1(42); // expected-error {{call to global function g1 not configured}}
 
   t1(1);