[PATCH] D15858: Warn undeclared identifiers in CUDA kernel calls

[Jason Henline via cfe-commits]

jhen created this revision.
jhen added reviewers: tra, jlebar.
jhen added a subscriber: cfe-commits.

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

http://reviews.llvm.org/D15858

Files:
  include/clang/AST/ExprCXX.h
  test/SemaCUDA/kernel-call.cu

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'}}
 }
Index: include/clang/AST/ExprCXX.h
===================================================================
--- include/clang/AST/ExprCXX.h
+++ include/clang/AST/ExprCXX.h
@@ -171,7 +171,15 @@
     return cast_or_null<CallExpr>(getPreArg(CONFIG));
   }
   CallExpr *getConfig() { return cast_or_null<CallExpr>(getPreArg(CONFIG)); }
-  void setConfig(CallExpr *E) { setPreArg(CONFIG, E); }
+  void setConfig(CallExpr *E) {
+    setPreArg(CONFIG, E);
+    setValueDependent(isValueDependent() || E->isValueDependent());
+    setTypeDependent(isTypeDependent() || E->isTypeDependent());
+    setInstantiationDependent(isInstantiationDependent() ||
+                              E->isInstantiationDependent());
+    setContainsUnexpandedParameterPack(containsUnexpandedParameterPack() ||
+                                       E->containsUnexpandedParameterPack());
+  }
 
   static bool classof(const Stmt *T) {
     return T->getStmtClass() == CUDAKernelCallExprClass;


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D15858.43897.patch
Type: text/x-patch
Size: 1347 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20160104/a8ecbec4/attachment.bin>