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

[Jason Henline via cfe-commits]

jhen marked an inline comment as done.

================
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:
> jhen wrote:
> > Thanks for bringing this up. While trying to find tests that dealt with each dependence individually, I came to realize that value and type dependence should not be set for the CUDAKernelCallExpr node because it's value is always void. So, I removed the propagation of those two dependencies.
> > 
> > Then, while looking for a test that could handle the parameter pack information, I realized that it was opening up a whole new can of worms and that the triple-angle-bracket syntax does not currently support variadic templates. I decided that parameter packs should be handled as a separate bug, so I removed them from this patch.
> > 
> > The instantiation dependence propagation is still valid, though, because it just represents whether a template parameter is present anywhere in the expression, so I left it in. Correctly tracking instantiation dependence in enough to fix the bug this patch was meant to fix, so I think it is the only change that should be made in this patch.
> What happens if an unexpanded pack is used within the kernel arguments of a CUDA kernel call? Do we already reject that? Are there tests for that somewhere?
There don't seem to be any tests currently that handle this case.

The case I had in mind for an unexpanded parameter pack was something like the following:

  __global__ void kernel() {}

  template <int ...Dimensions> kernel_wrapper() {
    kernel<<<Dimensions...>>>();
  }

This currently leads to a warning at the time of parsing that says the closing ">>>" is not found. I believe the cause is that the argument list is parsed as a simple argument list, so it doesn't handle the ellipsis correctly. I experimented with using standard (non-simple) parsing for the argument list, but that led to failures in other unit tests where ">>" wasn't being warned correctly in C++98 mode. I'm planning to file a bug for this (at least to fix the warning if not to allow the construction) and deal with it in a later patch. Does that sound reasonable?


http://reviews.llvm.org/D15858