Skip to content
This repository has been archived by the owner on Apr 23, 2020. It is now read-only.

Commit

Permalink
CUDA: add separate diagnostics for too few/many exec config args
Browse files Browse the repository at this point in the history 
git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@140977 91177308-0d34-0410-b5e6-96231b3b80d8
  • Loading branch information
@pcc
pcc committed Oct 2, 2011
1 parent af15b4d commit 1f24076
Show file tree
Hide file tree
Showing 4 changed files with 31 additions and 17 deletions.
12 changes: 8 additions & 4 deletions include/clang/Basic/DiagnosticSemaKinds.td
View file
Expand Up @@ -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">;
Expand Down
8 changes: 5 additions & 3 deletions include/clang/Sema/Sema.h
View file
Expand Up @@ -2454,19 +2454,21 @@ class Sema {
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,
Expand Down
26 changes: 16 additions & 10 deletions lib/Sema/SemaExpr.cpp
View file
Expand Up @@ -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)
Expand All @@ -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.
Expand All @@ -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;

Expand All @@ -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;

Expand Down Expand Up @@ -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.
Expand Down Expand Up @@ -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
Expand All @@ -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.
Expand Down Expand Up @@ -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.
Expand Down Expand Up @@ -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!");
Expand Down
2 changes: 2 additions & 0 deletions test/SemaCUDA/kernel-call.cu
View file
Expand Up @@ -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);

Expand Down

0 comments on commit 1f24076

Please sign in to comment.