[PATCH] D40275: [CUDA] Report "unsupported VLA" errors only on device side.

[Artem Belevich via Phabricator via cfe-commits]

tra added inline comments.


================
Comment at: clang/lib/Sema/SemaType.cpp:2188
+               !Context.getTargetInfo().isVLASupported() &&
+               shouldDiagnoseTargetSupportFromOpenMP()) {
+      // Some targets don't support VLAs.
----------------
tra wrote:
> rjmccall wrote:
> > Please write this check so that it trips in an "ordinary" build on a target that just happens to not support VLAs, something like:
> > 
> >   else if (!Context.getTargetInfo().isVLASupported() && shouldDiagnoseTargetSupportFromOpenMP())
> > 
> > If you want to include the explicit OpenMP check there, it would need to be:
> > 
> >   else if (!Context.getTargetInfo().isVLASupported() && (!getLangOpts().OpenMP || shouldDiagnoseTargetSupportFromOpenMP()))
> > 
> > but I think the first looks better.
> > 
> > The CUDA and OpenMP paths here seem to be trying to achieve analogous things; it's unfortunate that we can't find a way to unify their approaches, even if we'd eventually want to use different diagnostic text.  I imagine that the target-environment language restrictions are basically the same, since they arise for the same fundamental reasons, so all the places using CUDADiagIfDeviceCode are likely to have a check for shouldDiagnoseTargetSupportFromOpenMP() and vice-versa.
> The problem is that in CUDA we can't just do 
> ```
> if (!Context.getTargetInfo().isVLASupported() && shouldDiagnoseTargetSupportFromOpenMP())
>    Diag(Loc, diag::err_vla_unsupported);
> ```
> 
> In some situations diag messages will only be emitted if we attempt to generate unsupported code on device side.
> Consider 
> ```
> __host__ __device__ void foo(int n) {
>   int vla[n];
> }
> ```
> 
> When Sema sees this code during compilation, it can not tell whether there is an error. Calling foo from the host code is perfectly valid. Calling it from device code is not. `CUDADiagIfDeviceCode` creates 'postponed' diagnostics which only gets emitted if we ever need to generate code for the function on device.
> 
> So, while CUDA and OpenMP do similar things, they are not quite the same.  If your goal to generalize CUDA and OpenMP handling, then it would have to be folded into diagnostic-emitting itself and we'll need an analog of CUDADiagIfDeviceCode which can handle both OpenMP and CUDA. 
> E.g. something like this:
> ```
> ??? DiagIfDeviceCode(???) {
>    if (OpenCL || (OpenMP && shouldDiagnoseTargetSupportFromOpenMP()))
>        Diag(...);
>    else if (CUDA)
>        CUDADiagIfDeviceCode()
> } 
> 
> ...
> 
> if (!Context.getTargetInfo().isVLASupported()) 
>    DiagIfDeviceCode();
> 
> ```
> 
> Would that work for you?
There's another issue with this approach -- diagnostics itself. Each dialect has its own. Specifically CUDA diags have details that are relevant only to CUDA. I suspect OpenMP has something specific as well. If we insist emitting only one kind of error for particular case across all dialects, we'll have to stick to bare bones "feature X is not supported" which will not have sufficient details to explain why the error was triggered in CUDA.

IMO dialect-specific handling of cuda errors in this case is the lesser evil. 

I'll update the patch to handle non-cuda cases the way  you suggested.


https://reviews.llvm.org/D40275