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;
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;
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
