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

Jason Henline via cfe-commits cfe-commits at lists.llvm.org
Wed Jan 6 14:56:44 PST 2016


jhen added inline comments.

================
Comment at: include/clang/AST/ExprCXX.h:181
@@ +180,3 @@
+    assert(
+        !IsConfigSet &&
+        "CUDAKernelCallExpr.setConfig can only be called once per instance.");
----------------
rsmith wrote:
> Perhaps `assert(!getPreArg(CONFIG))` instead of storing a separate flag?
Yes, I think that is much nicer. But I am worried that the config pre-arg might not be initialized before I check it, so in my next revision of this patch I've also added code in the CallExpr constructors to make sure that pre-args are always initialized to zero.

Unfortunately, the simple way to do this leads to the entire array of function argument pointers being initialized for every function call expression construction (not just CUDA calls). Do you think I should try to avoid this overhead by only initializing the pre-arguments (either in the CallExpr code for everyone to share, or just in the CUDA code)?

================
Comment at: test/SemaCUDA/kernel-call.cu:27
@@ -26,1 +26,3 @@
+
+  g1<<<undeclared, 1>>>(42); // expected-error {{use of undeclared identifier 'undeclared'}}
 }
----------------
rsmith wrote:
> Your approach for that testcase seems fine, but it's not a test for the right thing as it doesn't have an unexpanded pack within the kernel call args. Here's a testcase for that scenario:
> 
>   template<int ...Dimensions> void kernel_wrapper() {
>     void (*fs[])() = {
>       []{ kernel<<<Dimensions>>>(); } ...
>     };
>   }
Excellent! Thank you for that test case. I suspected that I wasn't really exercising the "unexpanded" parameter pack with my code.


http://reviews.llvm.org/D15858





More information about the cfe-commits mailing list