Index: include/clang/AST/ExprCXX.h
===================================================================
--- include/clang/AST/ExprCXX.h
+++ include/clang/AST/ExprCXX.h
@@ -171,7 +171,18 @@
     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 can only be called once per class instance.
+  void setConfig(CallExpr *E) {
+    assert(
+        !getConfig() &&
+        "CUDAKernelCallExpr.setConfig can only be called once per instance.");
+    setPreArg(CONFIG, E);
+    setInstantiationDependent(isInstantiationDependent() ||
+                              E->isInstantiationDependent());
+  }
 
   static bool classof(const Stmt *T) {
     return T->getStmtClass() == CUDAKernelCallExprClass;
Index: lib/AST/Expr.cpp
===================================================================
--- lib/AST/Expr.cpp
+++ lib/AST/Expr.cpp
@@ -1148,7 +1148,7 @@
          fn->containsUnexpandedParameterPack()),
     NumArgs(args.size()) {
 
-  SubExprs = new (C) Stmt*[args.size()+PREARGS_START+NumPreArgs];
+  SubExprs = new (C) Stmt*[args.size()+PREARGS_START+NumPreArgs]();
   SubExprs[FN] = fn;
   for (unsigned i = 0; i != args.size(); ++i) {
     if (args[i]->isTypeDependent())
@@ -1179,7 +1179,7 @@
                    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;
 }
 
Index: test/SemaCUDA/kernel-call.cu
===================================================================
--- test/SemaCUDA/kernel-call.cu
+++ test/SemaCUDA/kernel-call.cu
@@ -23,4 +23,6 @@
 
   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'}}
 }