diff options
| author | Peter Collingbourne <peter@pcc.me.uk> | 2011-10-02 23:49:15 +0000 |
|---|---|---|
| committer | Peter Collingbourne <peter@pcc.me.uk> | 2011-10-02 23:49:15 +0000 |
| commit | 34a20b081e43ce3d7c6752364a2358ae00050494 (patch) | |
| tree | 6f443d8530e9fa42c3c3b279897359476785b5f1 | |
| parent | 7aea69d9490dc6a07ed3165982ff38a7a8077cb2 (diff) | |
| download | bcm5719-llvm-34a20b081e43ce3d7c6752364a2358ae00050494.tar.gz bcm5719-llvm-34a20b081e43ce3d7c6752364a2358ae00050494.zip | |
CUDA: diagnose unconfigured calls to global functions
llvm-svn: 140975
| -rw-r--r-- | clang/include/clang/Basic/DiagnosticSemaKinds.td | 2 | ||||
| -rw-r--r-- | clang/lib/Sema/SemaExpr.cpp | 5 | ||||
| -rw-r--r-- | clang/test/SemaCUDA/kernel-call.cu | 1 |
3 files changed, 8 insertions, 0 deletions
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 14b36e45985..357f1af6114 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/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/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index d0a3f6b5ecc..eb5678e7f00 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/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/clang/test/SemaCUDA/kernel-call.cu b/clang/test/SemaCUDA/kernel-call.cu index 7bc7ae11315..9a22aa77ef4 100644 --- a/clang/test/SemaCUDA/kernel-call.cu +++ b/clang/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); |
