From 02f9ccaed645dc8d70e726457f43802469d4938d Mon Sep 17 00:00:00 2001 From: Justin Lebar Date: Thu, 14 Jan 2016 23:31:30 +0000 Subject: [PATCH] [CUDA] Warn undeclared identifiers in CUDA kernel calls Value, type, and instantiation dependence were not being handled correctly for CUDAKernelCallExpr AST nodes. As a result, if an undeclared identifier was used in the triple-angle-bracket kernel call configuration, there would be no error during parsing, and there would be a crash during code gen. This patch makes sure that an error will be issued during parsing in this case, just as there would be for any other use of an undeclared identifier in C++. Patch by Jason Henline. Reviewers: jlebar, rsmith Differential Revision: http://reviews.llvm.org/D15858 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@257839 91177308-0d34-0410-b5e6-96231b3b80d8 --- include/clang/AST/Expr.h | 10 ++++-- include/clang/AST/ExprCXX.h | 26 ++++++++++----- lib/AST/Expr.cpp | 53 ++++++++++++++++++------------ test/SemaCUDA/cxx11-kernel-call.cu | 10 ++++++ test/SemaCUDA/kernel-call.cu | 2 ++ 5 files changed, 69 insertions(+), 32 deletions(-) create mode 100644 test/SemaCUDA/cxx11-kernel-call.cu diff --git a/include/clang/AST/Expr.h b/include/clang/AST/Expr.h index 38733eee82..bc5dab754c 100644 --- a/include/clang/AST/Expr.h +++ b/include/clang/AST/Expr.h @@ -2137,11 +2137,15 @@ class CallExpr : public Expr { unsigned NumArgs; SourceLocation RParenLoc; + void updateDependenciesFromArg(Expr *Arg); + protected: // These versions of the constructor are for derived classes. - CallExpr(const ASTContext& C, StmtClass SC, Expr *fn, unsigned NumPreArgs, - ArrayRef args, QualType t, ExprValueKind VK, - SourceLocation rparenloc); + CallExpr(const ASTContext &C, StmtClass SC, Expr *fn, + ArrayRef preargs, ArrayRef args, QualType t, + ExprValueKind VK, SourceLocation rparenloc); + CallExpr(const ASTContext &C, StmtClass SC, Expr *fn, ArrayRef args, + QualType t, ExprValueKind VK, SourceLocation rparenloc); CallExpr(const ASTContext &C, StmtClass SC, unsigned NumPreArgs, EmptyShell Empty); diff --git a/include/clang/AST/ExprCXX.h b/include/clang/AST/ExprCXX.h index 6821274986..2b8c0ea953 100644 --- a/include/clang/AST/ExprCXX.h +++ b/include/clang/AST/ExprCXX.h @@ -66,8 +66,7 @@ public: CXXOperatorCallExpr(ASTContext& C, OverloadedOperatorKind Op, Expr *fn, ArrayRef args, QualType t, ExprValueKind VK, SourceLocation operatorloc, bool fpContractable) - : CallExpr(C, CXXOperatorCallExprClass, fn, 0, args, t, VK, - operatorloc), + : CallExpr(C, CXXOperatorCallExprClass, fn, args, t, VK, operatorloc), Operator(Op), FPContractable(fpContractable) { Range = getSourceRangeImpl(); } @@ -125,7 +124,7 @@ class CXXMemberCallExpr : public CallExpr { public: CXXMemberCallExpr(ASTContext &C, Expr *fn, ArrayRef args, QualType t, ExprValueKind VK, SourceLocation RP) - : CallExpr(C, CXXMemberCallExprClass, fn, 0, args, t, VK, RP) {} + : CallExpr(C, CXXMemberCallExprClass, fn, args, t, VK, RP) {} CXXMemberCallExpr(ASTContext &C, EmptyShell Empty) : CallExpr(C, CXXMemberCallExprClass, Empty) { } @@ -160,9 +159,7 @@ public: CUDAKernelCallExpr(ASTContext &C, Expr *fn, CallExpr *Config, ArrayRef args, QualType t, ExprValueKind VK, SourceLocation RP) - : CallExpr(C, CUDAKernelCallExprClass, fn, END_PREARG, args, t, VK, RP) { - setConfig(Config); - } + : CallExpr(C, CUDAKernelCallExprClass, fn, Config, args, t, VK, RP) {} CUDAKernelCallExpr(ASTContext &C, EmptyShell Empty) : CallExpr(C, CUDAKernelCallExprClass, END_PREARG, Empty) { } @@ -171,7 +168,20 @@ public: return cast_or_null(getPreArg(CONFIG)); } CallExpr *getConfig() { return cast_or_null(getPreArg(CONFIG)); } - void setConfig(CallExpr *E) { setPreArg(CONFIG, E); } + + /// \brief Sets the kernel configuration expression. + /// + /// Note that this method cannot be called if config has already been set to a + /// non-null value. + void setConfig(CallExpr *E) { + assert(!getConfig() && + "Cannot call setConfig if config is not null"); + setPreArg(CONFIG, E); + setInstantiationDependent(isInstantiationDependent() || + E->isInstantiationDependent()); + setContainsUnexpandedParameterPack(containsUnexpandedParameterPack() || + E->containsUnexpandedParameterPack()); + } static bool classof(const Stmt *T) { return T->getStmtClass() == CUDAKernelCallExprClass; @@ -398,7 +408,7 @@ public: UserDefinedLiteral(const ASTContext &C, Expr *Fn, ArrayRef Args, QualType T, ExprValueKind VK, SourceLocation LitEndLoc, SourceLocation SuffixLoc) - : CallExpr(C, UserDefinedLiteralClass, Fn, 0, Args, T, VK, LitEndLoc), + : CallExpr(C, UserDefinedLiteralClass, Fn, Args, T, VK, LitEndLoc), UDSuffixLoc(SuffixLoc) {} explicit UserDefinedLiteral(const ASTContext &C, EmptyShell Empty) : CallExpr(C, UserDefinedLiteralClass, Empty) {} diff --git a/lib/AST/Expr.cpp b/lib/AST/Expr.cpp index 52f34df435..42d7ed3ac6 100644 --- a/lib/AST/Expr.cpp +++ b/lib/AST/Expr.cpp @@ -1138,28 +1138,23 @@ OverloadedOperatorKind UnaryOperator::getOverloadedOperator(Opcode Opc) { // Postfix Operators. //===----------------------------------------------------------------------===// -CallExpr::CallExpr(const ASTContext& C, StmtClass SC, Expr *fn, - unsigned NumPreArgs, ArrayRef args, QualType t, +CallExpr::CallExpr(const ASTContext &C, StmtClass SC, Expr *fn, + ArrayRef preargs, ArrayRef args, QualType t, ExprValueKind VK, SourceLocation rparenloc) - : Expr(SC, t, VK, OK_Ordinary, - fn->isTypeDependent(), - fn->isValueDependent(), - fn->isInstantiationDependent(), - fn->containsUnexpandedParameterPack()), - NumArgs(args.size()) { - - SubExprs = new (C) Stmt*[args.size()+PREARGS_START+NumPreArgs]; + : Expr(SC, t, VK, OK_Ordinary, fn->isTypeDependent(), + fn->isValueDependent(), fn->isInstantiationDependent(), + fn->containsUnexpandedParameterPack()), + NumArgs(args.size()) { + + unsigned NumPreArgs = preargs.size(); + SubExprs = new (C) Stmt *[args.size()+PREARGS_START+NumPreArgs]; SubExprs[FN] = fn; + for (unsigned i = 0; i != NumPreArgs; ++i) { + updateDependenciesFromArg(preargs[i]); + SubExprs[i+PREARGS_START] = preargs[i]; + } for (unsigned i = 0; i != args.size(); ++i) { - if (args[i]->isTypeDependent()) - ExprBits.TypeDependent = true; - if (args[i]->isValueDependent()) - ExprBits.ValueDependent = true; - if (args[i]->isInstantiationDependent()) - ExprBits.InstantiationDependent = true; - if (args[i]->containsUnexpandedParameterPack()) - ExprBits.ContainsUnexpandedParameterPack = true; - + updateDependenciesFromArg(args[i]); SubExprs[i+PREARGS_START+NumPreArgs] = args[i]; } @@ -1167,9 +1162,14 @@ CallExpr::CallExpr(const ASTContext& C, StmtClass SC, Expr *fn, RParenLoc = rparenloc; } +CallExpr::CallExpr(const ASTContext &C, StmtClass SC, Expr *fn, + ArrayRef args, QualType t, ExprValueKind VK, + SourceLocation rparenloc) + : CallExpr(C, SC, fn, ArrayRef(), args, t, VK, rparenloc) {} + CallExpr::CallExpr(const ASTContext &C, Expr *fn, ArrayRef args, QualType t, ExprValueKind VK, SourceLocation rparenloc) - : CallExpr(C, CallExprClass, fn, /*NumPreArgs=*/0, args, t, VK, rparenloc) { + : CallExpr(C, CallExprClass, fn, ArrayRef(), args, t, VK, rparenloc) { } CallExpr::CallExpr(const ASTContext &C, StmtClass SC, EmptyShell Empty) @@ -1179,10 +1179,21 @@ CallExpr::CallExpr(const ASTContext &C, StmtClass SC, unsigned NumPreArgs, EmptyShell Empty) : Expr(SC, Empty), SubExprs(nullptr), NumArgs(0) { // FIXME: Why do we allocate this? - SubExprs = new (C) Stmt*[PREARGS_START+NumPreArgs]; + SubExprs = new (C) Stmt*[PREARGS_START+NumPreArgs](); CallExprBits.NumPreArgs = NumPreArgs; } +void CallExpr::updateDependenciesFromArg(Expr *Arg) { + if (Arg->isTypeDependent()) + ExprBits.TypeDependent = true; + if (Arg->isValueDependent()) + ExprBits.ValueDependent = true; + if (Arg->isInstantiationDependent()) + ExprBits.InstantiationDependent = true; + if (Arg->containsUnexpandedParameterPack()) + ExprBits.ContainsUnexpandedParameterPack = true; +} + Decl *CallExpr::getCalleeDecl() { Expr *CEE = getCallee()->IgnoreParenImpCasts(); diff --git a/test/SemaCUDA/cxx11-kernel-call.cu b/test/SemaCUDA/cxx11-kernel-call.cu new file mode 100644 index 0000000000..0d80125633 --- /dev/null +++ b/test/SemaCUDA/cxx11-kernel-call.cu @@ -0,0 +1,10 @@ +// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify %s + +#include "Inputs/cuda.h" + +__global__ void k1() {} + +template void k1Wrapper() { + void (*f)() = [] { k1<<>>(); }; // expected-error {{initializer contains unexpanded parameter pack 'Dimensions'}} + void (*g[])() = { [] { k1<<>>(); } ... }; // ok +} diff --git a/test/SemaCUDA/kernel-call.cu b/test/SemaCUDA/kernel-call.cu index 9a3d86c47f..47d7762f59 100644 --- a/test/SemaCUDA/kernel-call.cu +++ b/test/SemaCUDA/kernel-call.cu @@ -23,4 +23,6 @@ int main(void) { int (*fp)(int) = h2; fp<<<1, 1>>>(42); // expected-error {{must have void return type}} + + g1<<>>(42); // expected-error {{use of undeclared identifier 'undeclared'}} } -- 2.50.1