[cfe-commits] r140977 - in /cfe/trunk: include/clang/Basic/DiagnosticSemaKinds.td include/clang/Sema/Sema.h lib/Sema/SemaExpr.cpp test/SemaCUDA/kernel-call.cu
Peter Collingbourne
peter at pcc.me.uk
Sun Oct 2 16:49:30 PDT 2011
Author: pcc
Date: Sun Oct 2 18:49:29 2011
New Revision: 140977
URL: http://llvm.org/viewvc/llvm-project?rev=140977&view=rev
Log:
CUDA: add separate diagnostics for too few/many exec config args
Modified:
cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
cfe/trunk/include/clang/Sema/Sema.h
cfe/trunk/lib/Sema/SemaExpr.cpp
cfe/trunk/test/SemaCUDA/kernel-call.cu
Modified: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td?rev=140977&r1=140976&r2=140977&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td (original)
+++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td Sun Oct 2 18:49:29 2011
@@ -3954,16 +3954,20 @@
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">;
Modified: cfe/trunk/include/clang/Sema/Sema.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Sema/Sema.h?rev=140977&r1=140976&r2=140977&view=diff
==============================================================================
--- cfe/trunk/include/clang/Sema/Sema.h (original)
+++ cfe/trunk/include/clang/Sema/Sema.h Sun Oct 2 18:49:29 2011
@@ -2454,19 +2454,21 @@
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,
Modified: cfe/trunk/lib/Sema/SemaExpr.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaExpr.cpp?rev=140977&r1=140976&r2=140977&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaExpr.cpp (original)
+++ cfe/trunk/lib/Sema/SemaExpr.cpp Sun Oct 2 18:49:29 2011
@@ -3259,7 +3259,8 @@
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 @@
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 @@
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 @@
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 @@
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 @@
NDecl = cast<MemberExpr>(NakedFn)->getMemberDecl();
return BuildResolvedCallExpr(Fn, NDecl, LParenLoc, Args, NumArgs, RParenLoc,
- ExecConfig);
+ ExecConfig, IsExecConfig);
}
ExprResult
@@ -3555,7 +3560,8 @@
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 @@
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 @@
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!");
Modified: cfe/trunk/test/SemaCUDA/kernel-call.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/kernel-call.cu?rev=140977&r1=140976&r2=140977&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/kernel-call.cu (original)
+++ cfe/trunk/test/SemaCUDA/kernel-call.cu Sun Oct 2 18:49:29 2011
@@ -14,6 +14,8 @@
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);
More information about the cfe-commits
mailing list