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

[Jason Henline via cfe-commits]

jhen updated this revision to Diff 44131.
jhen added a reviewer: jlebar.
jhen removed a subscriber: jlebar.
jhen added a comment.

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


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,11 @@
     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);
+    setInstantiationDependent(isInstantiationDependent() ||
+                              E->isInstantiationDependent());
+  }
 
   static bool classof(const Stmt *T) {
     return T->getStmtClass() == CUDAKernelCallExprClass;


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D15858.44131.patch
Type: text/x-patch
Size: 1056 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20160106/19d238e2/attachment.bin>