From 8591a7f5215f6469603e6bc8f4fdbcc78e215ab9 Mon Sep 17 00:00:00 2001
From: Peter Collingbourne <peter@pcc.me.uk>
Date: Sun, 2 Oct 2011 23:49:15 +0000
Subject: [PATCH] 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 | 2 ++
lib/Sema/SemaExpr.cpp | 5 +++++
test/SemaCUDA/kernel-call.cu | 1 +
3 files changed, 8 insertions(+)
diff --git a/include/clang/Basic/DiagnosticSemaKinds.td b/include/clang/Basic/DiagnosticSemaKinds.td
index 14b36e45985..357f1af6114 100644
--- a/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/include/clang/Basic/DiagnosticSemaKinds.td
@@ -3986,6 +3986,8 @@ 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_global_call_not_config : Error<
+ "call to global function %0 not configured">;
def err_cannot_pass_objc_interface_to_vararg : Error<
diff --git a/lib/Sema/SemaExpr.cpp b/lib/Sema/SemaExpr.cpp
index d0a3f6b5ecc..eb5678e7f00 100644
--- a/lib/Sema/SemaExpr.cpp
+++ b/lib/Sema/SemaExpr.cpp
@@ -3655,6 +3655,11 @@ Sema::BuildResolvedCallExpr(Expr *Fn, NamedDecl *NDecl,
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());
}
}
diff --git a/test/SemaCUDA/kernel-call.cu b/test/SemaCUDA/kernel-call.cu
index 7bc7ae11315..9a22aa77ef4 100644
--- a/test/SemaCUDA/kernel-call.cu
+++ b/test/SemaCUDA/kernel-call.cu
@@ -13,6 +13,7 @@ int h2(int x) { return 1; }
int main(void) {
g1<<<1, 1>>>(42);
+ g1(42); // expected-error {{call to global function g1 not configured}}
t1(1);
--
GitLab