Warn undeclared identifiers in CUDA kernel calls

[jhen]

Summary

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++.

[jlebar]

test/SemaCUDA/kernel-call.cu

27 ↗

(On Diff #43897)

We set four things in setConfig -- does this test fail if any one of them is commented out? If not, is it hard to add additional tests that will cover the changes?

[jhen]

• Correct dependence info in CUDA kernel call AST
  
  This patch removes the propagation of type and value dependence and the propagation of information on unexpanded parameter packs from the CUDA kernel configuration function call expression to its parent CUDA kernel call expression AST node. It does, however, maintain the propagation of instantiation dependence between those nodes, as introduced in the previous revision of this patch.
  
  The last patch should not have propagated value and type dependence from the CUDA kernel config function to the entire CUDA kernel call expression AST node. The reason is that the CUDA kernel call expression has a void value, so it's value cannot depend on template types or values, it is always simply void.
  
  However, the CUDA kernel call expression node can contain template arguments, so it can be instantiation dependent. That means that the instantiation dependence should be propagated from the config call to the kernel call node. The instantiation dependence propagation is also sufficient to fix the crashing bug that results from using an undeclared identifier as a config argument.
  
  As for tracking unexpanded parameter packs, it is not yet clear how the CUDA triple-angle-bracket syntax will interoperate with variadic templates, so I will leave that propagation out of this patch and it can be dealt with later.

[jlebar]

Looks sane to me, although I have no idea what I'm doing here; you should probably get someone else's approval.

[jhen]

test/SemaCUDA/kernel-call.cu

27 ↗

(On Diff #44131)

Thanks for bringing this up. While trying to find tests that dealt with each dependence individually, I came to realize that value and type dependence should not be set for the CUDAKernelCallExpr node because it's value is always void. So, I removed the propagation of those two dependencies.

Then, while looking for a test that could handle the parameter pack information, I realized that it was opening up a whole new can of worms and that the triple-angle-bracket syntax does not currently support variadic templates. I decided that parameter packs should be handled as a separate bug, so I removed them from this patch.

The instantiation dependence propagation is still valid, though, because it just represents whether a template parameter is present anywhere in the expression, so I left it in. Correctly tracking instantiation dependence in enough to fix the bug this patch was meant to fix, so I think it is the only change that should be made in this patch.

[rsmith]

include/clang/AST/ExprCXX.h

175 ↗

(On Diff #44131)

Can you assert that the argument is only set once here? (If we set it to something instantiation-dependent and then to something that isn't instantiation-dependent, we'd compute the wrong instantiation-dependence flag.) This function should only be called by the normal constructor and by people who called the EmptyShell constructor.

test/SemaCUDA/kernel-call.cu

27 ↗

(On Diff #44131)

What happens if an unexpanded pack is used within the kernel arguments of a CUDA kernel call? Do we already reject that? Are there tests for that somewhere?

[jhen]

• Assert setConfig only called once

[jhen]

test/SemaCUDA/kernel-call.cu

27 ↗

(On Diff #44131)

There don't seem to be any tests currently that handle this case.

The case I had in mind for an unexpanded parameter pack was something like the following:

__global__ void kernel() {}

template <int ...Dimensions> kernel_wrapper() {
  kernel<<<Dimensions...>>>();
}

This currently leads to a warning at the time of parsing that says the closing ">>>" is not found. I believe the cause is that the argument list is parsed as a simple argument list, so it doesn't handle the ellipsis correctly. I experimented with using standard (non-simple) parsing for the argument list, but that led to failures in other unit tests where ">>" wasn't being warned correctly in C++98 mode. I'm planning to file a bug for this (at least to fix the warning if not to allow the construction) and deal with it in a later patch. Does that sound reasonable?

[rsmith]

include/clang/AST/ExprCXX.h

181 ↗

(On Diff #44143)

Perhaps assert(!getPreArg(CONFIG)) instead of storing a separate flag?

test/SemaCUDA/kernel-call.cu

27 ↗

(On Diff #44143)

Your approach for that testcase seems fine, but it's not a test for the right thing as it doesn't have an unexpanded pack within the kernel call args. Here's a testcase for that scenario:

template<int ...Dimensions> void kernel_wrapper() {
  void (*fs[])() = {
    []{ kernel<<<Dimensions>>>(); } ...
  };
}

[jhen]

include/clang/AST/ExprCXX.h

181 ↗

(On Diff #44143)

Yes, I think that is much nicer. But I am worried that the config pre-arg might not be initialized before I check it, so in my next revision of this patch I've also added code in the CallExpr constructors to make sure that pre-args are always initialized to zero.

Unfortunately, the simple way to do this leads to the entire array of function argument pointers being initialized for every function call expression construction (not just CUDA calls). Do you think I should try to avoid this overhead by only initializing the pre-arguments (either in the CallExpr code for everyone to share, or just in the CUDA code)?

test/SemaCUDA/kernel-call.cu

27 ↗

(On Diff #44143)

Excellent! Thank you for that test case. I suspected that I wasn't really exercising the "unexpanded" parameter pack with my code.

[jhen]

• Use config ptr itself rather than boolean flag

[rsmith]

include/clang/AST/ExprCXX.h

180 ↗

(On Diff #44167)

My preference would be to pass the CallExpr constructor an ArrayRef<Expr*> for the preargs and have it initialize them itself.

[jhen]

• CallExpr ctor for ArrayRef of pre-args

[jhen]

include/clang/AST/ExprCXX.h

177 ↗

(On Diff #44229)

Sounds good. The current version introduces a new constructor for CallExpr that takes an ArrayRef<Expr*> for the preargs, and the CUDAKernelCallExpr constructor now passes the Config argument explicitly to this constructor. This new constructor also replaces the old one that took the number of preargs. In the new constructor I designed the dependency handling for the preargs to mirror the way it is done for regular function arguments, and I extracted a little helper function so as not to repeat that code.

I also introduced another new CallExpr constructor to handle the common case of zero preargs, and replaced calls to the old NumPreArgs constructor where zero NumPreArgs was passed explicitly with calls to this new constructor.

[jhen]

• Handle unexpanded parameter packs
  
  The changes for instantiation dependence also fix a bug with unexpanded parameter packs, so add a unit test for unexpanded parameter packs as well.

[jhen]

rsmith, I think the patch is ready to be committed. Please take a look if you have a moment. Thanks for your help.

[rsmith]

LGTM, thanks!

test/SemaCUDA/cxx11-kernel-call.cu

8 ↗

(On Diff #44538)

Perhaps also add another test:

void (*g[])() = { [] { k1<<<Dimensions, Dimensions>>>(); } ... }; // ok

[jhen]

• Add extra test for OK parameter pack

[jhen]

Thanks for the review rsmith!

cfe/trunk/include/clang/AST/Expr.h

 /// a subclass for overloaded operator calls that use operator syntax, e.g.,
 /// "str1 + str2" to resolve to a function call.
 class CallExpr : public Expr {
   enum { FN=0, PREARGS_START=1 };
   Stmt **SubExprs;
   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);
 
   Stmt *getPreArg(unsigned i) {
     assert(i < getNumPreArgs() && "Prearg access out of range!");
     return SubExprs[PREARGS_START+i];
   }
   const Stmt *getPreArg(unsigned i) const {

cfe/trunk/include/clang/AST/ExprCXX.h

class CXXOperatorCallExpr : public CallExpr {
   // set to false.
   unsigned FPContractable : 1;
 
   SourceRange getSourceRangeImpl() const LLVM_READONLY;
 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();
   }
   explicit CXXOperatorCallExpr(ASTContext& C, EmptyShell Empty) :
     CallExpr(C, CXXOperatorCallExprClass, Empty) { }
 
 
   /// \brief Returns the kind of overloaded operator that this

 /// function. The callee in either case is a MemberExpr that contains
 /// both the object argument and the member function, while the
 /// arguments are the arguments within the parentheses (not including
 /// the object argument).
 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) { }
 
   /// \brief Retrieves the implicit object argument for the member call.
   ///
   /// For example, in "x.f(5)", this returns the sub-expression "x".
   Expr *getImplicitObjectArgument() const;

 class CUDAKernelCallExpr : public CallExpr {
 private:
   enum { CONFIG, END_PREARG };
 
 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) { }
 
   const CallExpr *getConfig() const {
     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;
   }
 };
 
 /// \brief Abstract class common to all of the C++ "named"/"keyword" casts.
 ///

 class UserDefinedLiteral : public CallExpr {
   /// \brief The location of a ud-suffix within the literal.
   SourceLocation UDSuffixLoc;
 
 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) {}
 
   /// The kind of literal operator which is invoked.
   enum LiteralOperatorKind {
     LOK_Raw,      ///< Raw form: operator "" X (const char *)
     LOK_Template, ///< Raw form: operator "" X<cs...> ()

cfe/trunk/lib/AST/Expr.cpp

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(),
+    : 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];
   }
 
   CallExprBits.NumPreArgs = NumPreArgs;
   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)
     : CallExpr(C, SC, /*NumPreArgs=*/0, Empty) {}
 
 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();
     
   while (SubstNonTypeTemplateParmExpr *NTTP
                                 = dyn_cast<SubstNonTypeTemplateParmExpr>(CEE)) {
     CEE = NTTP->getReplacement()->IgnoreParenCasts();
   }
   

cfe/trunk/test/SemaCUDA/cxx11-kernel-call.cu

+// 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
+}

cfe/trunk/test/SemaCUDA/kernel-call.cu

int main(void) {
   g1<<<1, 1, 0, 0, 0>>>(42); // expected-error {{too many execution configuration arguments to kernel function call}}
 
   t1(1);
 
   h1<<<1, 1>>>(42); // expected-error {{kernel call to non-global function h1}}
 
   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'}}
 }