diff options
| -rw-r--r-- | include/clang/Basic/DiagnosticSemaKinds.td | 2 | ||||
| -rw-r--r-- | lib/Sema/SemaExpr.cpp | 5 | ||||
| -rw-r--r-- | test/SemaCUDA/kernel-call.cu | 1 |
3 files changed, 8 insertions, 0 deletions
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<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 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); |