summaryrefslogtreecommitdiffstats
diff options
context:
space:
mode:
Diffstat
-rw-r--r--clang/include/clang/AST/Expr.h10
-rw-r--r--clang/include/clang/AST/ExprCXX.h26
-rw-r--r--clang/lib/AST/Expr.cpp53
-rw-r--r--clang/test/SemaCUDA/cxx11-kernel-call.cu10
-rw-r--r--clang/test/SemaCUDA/kernel-call.cu2
5 files changed, 69 insertions, 32 deletions
diff --git a/clang/include/clang/AST/Expr.h b/clang/include/clang/AST/Expr.h
index 38733eee82c..bc5dab754c6 100644
--- a/clang/include/clang/AST/Expr.h
+++ b/clang/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<Expr*> args, QualType t, ExprValueKind VK,
- SourceLocation rparenloc);
+ CallExpr(const ASTContext &C, StmtClass SC, Expr *fn,
+ ArrayRef<Expr *> preargs, ArrayRef<Expr *> args, QualType t,
+ ExprValueKind VK, SourceLocation rparenloc);
+ CallExpr(const ASTContext &C, StmtClass SC, Expr *fn, ArrayRef<Expr *> args,
+ QualType t, ExprValueKind VK, SourceLocation rparenloc);
CallExpr(const ASTContext &C, StmtClass SC, unsigned NumPreArgs,
EmptyShell Empty);
diff --git a/clang/include/clang/AST/ExprCXX.h b/clang/include/clang/AST/ExprCXX.h
index 68212749864..2b8c0ea953e 100644
--- a/clang/include/clang/AST/ExprCXX.h
+++ b/clang/include/clang/AST/ExprCXX.h
@@ -66,8 +66,7 @@ public:
CXXOperatorCallExpr(ASTContext& C, OverloadedOperatorKind Op, Expr *fn,
ArrayRef<Expr*> 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<Expr*> 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<Expr*> 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<CallExpr>(getPreArg(CONFIG));
}
CallExpr *getConfig() { return cast_or_null<CallExpr>(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<Expr*> 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/clang/lib/AST/Expr.cpp b/clang/lib/AST/Expr.cpp
index 52f34df4356..42d7ed3ac6f 100644
--- a/clang/lib/AST/Expr.cpp
+++ b/clang/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<Expr*> args, QualType t,
+CallExpr::CallExpr(const ASTContext &C, StmtClass SC, Expr *fn,
+ ArrayRef<Expr *> preargs, ArrayRef<Expr *> 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<Expr *> args, QualType t, ExprValueKind VK,
+ SourceLocation rparenloc)
+ : CallExpr(C, SC, fn, ArrayRef<Expr *>(), args, t, VK, rparenloc) {}
+
CallExpr::CallExpr(const ASTContext &C, Expr *fn, ArrayRef<Expr *> args,
QualType t, ExprValueKind VK, SourceLocation rparenloc)
- : CallExpr(C, CallExprClass, fn, /*NumPreArgs=*/0, args, t, VK, rparenloc) {
+ : CallExpr(C, CallExprClass, fn, ArrayRef<Expr *>(), 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/clang/test/SemaCUDA/cxx11-kernel-call.cu b/clang/test/SemaCUDA/cxx11-kernel-call.cu
new file mode 100644
index 00000000000..0d801256330
--- /dev/null
+++ b/clang/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<int ...Dimensions> void k1Wrapper() {
+ void (*f)() = [] { k1<<<Dimensions, Dimensions>>>(); }; // expected-error {{initializer contains unexpanded parameter pack 'Dimensions'}}
+ void (*g[])() = { [] { k1<<<Dimensions, Dimensions>>>(); } ... }; // ok
+}
diff --git a/clang/test/SemaCUDA/kernel-call.cu b/clang/test/SemaCUDA/kernel-call.cu
index 9a3d86c47fa..47d7762f59e 100644
--- a/clang/test/SemaCUDA/kernel-call.cu
+++ b/clang/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<<<undeclared, 1>>>(42); // expected-error {{use of undeclared identifier 'undeclared'}}
}
OpenPOWER on IntegriCloud