maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

CUDA: add separate diagnostics for too few/many exec config args 

llvm-svn: 140977
This commit is contained in:
Peter Collingbourne 2011-10-02 23:49:29 +00:00
parent 740afe276f
commit 619a8c7df3
4 changed files with 31 additions and 17 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

12
clang/include/clang/Basic/DiagnosticSemaKinds.td
View File

@ -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">;

8
clang/include/clang/Sema/Sema.h
View File

@ -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,

26
clang/lib/Sema/SemaExpr.cpp
View File

@ -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<MemberExpr>(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<FunctionDecl>(NDecl);
// Promote the function operand.
@ -3680,7 +3686,7 @@ Sema::BuildResolvedCallExpr(Expr *Fn, NamedDecl *NDecl,
if (const FunctionProtoType *Proto = dyn_cast<FunctionProtoType>(FuncT)) {
if (ConvertArgumentsForCall(TheCall, Fn, FDecl, Proto, Args, NumArgs,
RParenLoc))
RParenLoc, IsExecConfig))
return ExprError();
} else {
assert(isa<FunctionNoProtoType>(FuncT) && "Unknown FunctionType!");

2
clang/test/SemaCUDA/kernel-call.cu
View File

@ -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);