rjmccall 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:
> > tra wrote:
> > > 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.
> > If there really is interesting language-specific information to provide in
> > a diagnostic, I agree that it's hard to avoid having different code for
> > different targets. On the other hand, the CUDA-specific information in
> > this specific diagnostic seems unnecessary — does the user really care
> > about whether the function was '__device__' vs. '__host__ __device__'? in
> > fact, isn't the latter a bit misleading? — and I feel like a generic
> > "cannot use variable-length arrays when compiling for device 'nvptx64'"
> > would be perfectly satisfactory in both CUDA and OpenMP, and that's
> > probably true for almost all of these diagnostics.
> >
> > On a completely different note, I do want to point out that the only thing
> > you actually *need* to ban here is declaring a local variable of VLA type.
> > There's no reason at all to ban VLA types in general; they just compile to
> > extra arithmetic.
> Agreed. While host/device attributes are part of a function signature in
> CUDA, in this case only 'device' part makes the difference and therefore
> common-style reporting the way you suggest would be sufficient to indicate
> the error in the device-side code.
>
> As for the VLA types, clang can currently compile code with VLA arguments:
> https://godbolt.org/g/43hVu9
> Clang explicitly does not support GCC's VLA-in-structure extension.
> Are there any other use cases you can think of?
>
You can just make a VLA type locally (as a typedef, or just as a cast operand)
and use it for pointer arithmetic.
Repository:
rL LLVM
https://reviews.llvm.org/D40275
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
