diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 357f1af61149..7c6d23902a41 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -3954,16 +3954,20 @@ def note_function_with_incomplete_return_type_declared_here : Note< def err_call_incomplete_argument : Error< "argument type %0 is incomplete">; def err_typecheck_call_too_few_args : Error< - "too few arguments to %select{function|block|method}0 call, " + "too few %select{|||execution configuration }0arguments to " + "%select{function|block|method|kernel function}0 call, " "expected %1, have %2">; def err_typecheck_call_too_few_args_at_least : Error< - "too few arguments to %select{function|block|method}0 call, " + "too few %select{|||execution configuration }0arguments to " + "%select{function|block|method|kernel function}0 call, " "expected at least %1, have %2">; def err_typecheck_call_too_many_args : Error< - "too many arguments to %select{function|block|method}0 call, " + "too many %select{|||execution configuration }0arguments to " + "%select{function|block|method|kernel function}0 call, " "expected %1, have %2">; def err_typecheck_call_too_many_args_at_most : Error< - "too many arguments to %select{function|block|method}0 call, " + "too many %select{|||execution configuration }0arguments to " + "%select{function|block|method|kernel function}0 call, " "expected at most %1, have %2">; def note_callee_decl : Note< "%0 declared here">; diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 684bfc969b74..bf54bdc575ee 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -2454,19 +2454,21 @@ public: FunctionDecl *FDecl, const FunctionProtoType *Proto, Expr **Args, unsigned NumArgs, - SourceLocation RParenLoc); + SourceLocation RParenLoc, + bool ExecConfig = false); /// ActOnCallExpr - Handle a call to Fn with the specified array of arguments. /// This provides the location of the left/right parens and a list of comma /// locations. ExprResult ActOnCallExpr(Scope *S, Expr *Fn, SourceLocation LParenLoc, MultiExprArg ArgExprs, SourceLocation RParenLoc, - Expr *ExecConfig = 0); + Expr *ExecConfig = 0, bool IsExecConfig = false); ExprResult BuildResolvedCallExpr(Expr *Fn, NamedDecl *NDecl, SourceLocation LParenLoc, Expr **Args, unsigned NumArgs, SourceLocation RParenLoc, - Expr *Config = 0); + Expr *Config = 0, + bool IsExecConfig = false); ExprResult ActOnCUDAExecConfigExpr(Scope *S, SourceLocation LLLLoc, MultiExprArg ExecConfig, diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index 0d4c7bea171a..46d9fe64e0f3 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -3259,7 +3259,8 @@ Sema::ConvertArgumentsForCall(CallExpr *Call, Expr *Fn, FunctionDecl *FDecl, const FunctionProtoType *Proto, Expr **Args, unsigned NumArgs, - SourceLocation RParenLoc) { + SourceLocation RParenLoc, + bool IsExecConfig) { // Bail out early if calling a builtin with custom typechecking. // We don't need to do this in the if (FDecl) @@ -3272,6 +3273,10 @@ Sema::ConvertArgumentsForCall(CallExpr *Call, Expr *Fn, unsigned NumArgsInProto = Proto->getNumArgs(); bool Invalid = false; unsigned MinArgs = FDecl ? FDecl->getMinRequiredArguments() : NumArgsInProto; + unsigned FnKind = Fn->getType()->isBlockPointerType() + ? 1 /* block */ + : (IsExecConfig ? 3 /* kernel function (exec config) */ + : 0 /* function */); // If too few arguments are available (and we don't have default // arguments for the remaining parameters), don't make the call. @@ -3280,11 +3285,11 @@ Sema::ConvertArgumentsForCall(CallExpr *Call, Expr *Fn, Diag(RParenLoc, MinArgs == NumArgsInProto ? diag::err_typecheck_call_too_few_args : diag::err_typecheck_call_too_few_args_at_least) - << Fn->getType()->isBlockPointerType() + << FnKind << MinArgs << NumArgs << Fn->getSourceRange(); // Emit the location of the prototype. - if (FDecl && !FDecl->getBuiltinID()) + if (FDecl && !FDecl->getBuiltinID() && !IsExecConfig) Diag(FDecl->getLocStart(), diag::note_callee_decl) << FDecl; @@ -3301,13 +3306,13 @@ Sema::ConvertArgumentsForCall(CallExpr *Call, Expr *Fn, MinArgs == NumArgsInProto ? diag::err_typecheck_call_too_many_args : diag::err_typecheck_call_too_many_args_at_most) - << Fn->getType()->isBlockPointerType() + << FnKind << NumArgsInProto << NumArgs << Fn->getSourceRange() << SourceRange(Args[NumArgsInProto]->getLocStart(), Args[NumArgs-1]->getLocEnd()); // Emit the location of the prototype. - if (FDecl && !FDecl->getBuiltinID()) + if (FDecl && !FDecl->getBuiltinID() && !IsExecConfig) Diag(FDecl->getLocStart(), diag::note_callee_decl) << FDecl; @@ -3441,7 +3446,7 @@ static ExprResult rebuildUnknownAnyFunction(Sema &S, Expr *fn); ExprResult Sema::ActOnCallExpr(Scope *S, Expr *Fn, SourceLocation LParenLoc, MultiExprArg ArgExprs, SourceLocation RParenLoc, - Expr *ExecConfig) { + Expr *ExecConfig, bool IsExecConfig) { unsigned NumArgs = ArgExprs.size(); // Since this might be a postfix expression, get rid of ParenListExprs. @@ -3540,7 +3545,7 @@ Sema::ActOnCallExpr(Scope *S, Expr *Fn, SourceLocation LParenLoc, NDecl = cast(NakedFn)->getMemberDecl(); return BuildResolvedCallExpr(Fn, NDecl, LParenLoc, Args, NumArgs, RParenLoc, - ExecConfig); + ExecConfig, IsExecConfig); } ExprResult @@ -3555,7 +3560,8 @@ Sema::ActOnCUDAExecConfigExpr(Scope *S, SourceLocation LLLLoc, DeclRefExpr *ConfigDR = new (Context) DeclRefExpr( ConfigDecl, ConfigQTy, VK_LValue, LLLLoc); - return ActOnCallExpr(S, ConfigDR, LLLLoc, ExecConfig, GGGLoc, 0); + return ActOnCallExpr(S, ConfigDR, LLLLoc, ExecConfig, GGGLoc, 0, + /*IsExecConfig=*/true); } /// ActOnAsTypeExpr - create a new asType (bitcast) from the arguments. @@ -3590,7 +3596,7 @@ Sema::BuildResolvedCallExpr(Expr *Fn, NamedDecl *NDecl, SourceLocation LParenLoc, Expr **Args, unsigned NumArgs, SourceLocation RParenLoc, - Expr *Config) { + Expr *Config, bool IsExecConfig) { FunctionDecl *FDecl = dyn_cast_or_null(NDecl); // Promote the function operand. @@ -3680,7 +3686,7 @@ Sema::BuildResolvedCallExpr(Expr *Fn, NamedDecl *NDecl, if (const FunctionProtoType *Proto = dyn_cast(FuncT)) { if (ConvertArgumentsForCall(TheCall, Fn, FDecl, Proto, Args, NumArgs, - RParenLoc)) + RParenLoc, IsExecConfig)) return ExprError(); } else { assert(isa(FuncT) && "Unknown FunctionType!"); diff --git a/clang/test/SemaCUDA/kernel-call.cu b/clang/test/SemaCUDA/kernel-call.cu index 9a22aa77ef41..91b1d49e2d06 100644 --- a/clang/test/SemaCUDA/kernel-call.cu +++ b/clang/test/SemaCUDA/kernel-call.cu @@ -14,6 +14,8 @@ int h2(int x) { return 1; } int main(void) { g1<<<1, 1>>>(42); g1(42); // expected-error {{call to global function g1 not configured}} + g1<<<1>>>(42); // expected-error {{too few execution configuration arguments to kernel function call}} + g1<<<1, 1, 0, 0, 0>>>(42); // expected-error {{too many execution configuration arguments to kernel function call}} t1(1);