From 8591a7f5215f6469603e6bc8f4fdbcc78e215ab9 Mon Sep 17 00:00:00 2001 From: Peter Collingbourne 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 14b36e4598..357f1af611 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 d0a3f6b5ec..eb5678e7f0 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()) + 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 7bc7ae1131..9a22aa77ef 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); -- 2.40.0