[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;



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;

_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits