r257839 - [CUDA] Warn undeclared identifiers in CUDA kernel calls
Justin Lebar via cfe-commits
cfe-commits at lists.llvm.org
Thu Jan 14 15:31:34 PST 2016
Author: jlebar
Date: Thu Jan 14 17:31:30 2016
New Revision: 257839
URL: http://llvm.org/viewvc/llvm-project?rev=257839&view=rev
Log:
[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
Added:
cfe/trunk/test/SemaCUDA/cxx11-kernel-call.cu
Modified:
cfe/trunk/include/clang/AST/Expr.h
cfe/trunk/include/clang/AST/ExprCXX.h
cfe/trunk/lib/AST/Expr.cpp
cfe/trunk/test/SemaCUDA/kernel-call.cu
Modified: cfe/trunk/include/clang/AST/Expr.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/AST/Expr.h?rev=257839&r1=257838&r2=257839&view=diff
==============================================================================
--- cfe/trunk/include/clang/AST/Expr.h (original)
+++ cfe/trunk/include/clang/AST/Expr.h Thu Jan 14 17:31:30 2016
@@ -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);
Modified: cfe/trunk/include/clang/AST/ExprCXX.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/AST/ExprCXX.h?rev=257839&r1=257838&r2=257839&view=diff
==============================================================================
--- cfe/trunk/include/clang/AST/ExprCXX.h (original)
+++ cfe/trunk/include/clang/AST/ExprCXX.h Thu Jan 14 17:31:30 2016
@@ -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 CallExp
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) {}
Modified: cfe/trunk/lib/AST/Expr.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/AST/Expr.cpp?rev=257839&r1=257838&r2=257839&view=diff
==============================================================================
--- cfe/trunk/lib/AST/Expr.cpp (original)
+++ cfe/trunk/lib/AST/Expr.cpp Thu Jan 14 17:31:30 2016
@@ -1138,28 +1138,23 @@ OverloadedOperatorKind UnaryOperator::ge
// 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()) {
+ : 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];
+ 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,
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,
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();
Added: cfe/trunk/test/SemaCUDA/cxx11-kernel-call.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/cxx11-kernel-call.cu?rev=257839&view=auto
==============================================================================
--- cfe/trunk/test/SemaCUDA/cxx11-kernel-call.cu (added)
+++ cfe/trunk/test/SemaCUDA/cxx11-kernel-call.cu Thu Jan 14 17:31:30 2016
@@ -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
+}
Modified: cfe/trunk/test/SemaCUDA/kernel-call.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/kernel-call.cu?rev=257839&r1=257838&r2=257839&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/kernel-call.cu (original)
+++ cfe/trunk/test/SemaCUDA/kernel-call.cu Thu Jan 14 17:31:30 2016
@@ -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'}}
}
More information about the cfe-commits
mailing list