rjmccall added inline comments.
================
Comment at: include/clang/Basic/TargetInfo.h:944
+ /// \brief Whether target supports variable-length arrays.
+ bool isVLASupported() const { return VLASupported; }
+
----------------
ABataev wrote:
> ABataev wrote:
> > Hahnfeld wrote:
> > > rjmccall wrote:
> > > > Hahnfeld wrote:
> > > > > rjmccall wrote:
> > > > > > Hahnfeld wrote:
> > > > > > > rjmccall wrote:
> > > > > > > > The way you've written this makes it sound like "does the
> > > > > > > > target support VLAs?", but the actual semantic checks treat it
> > > > > > > > as "do OpenMP devices on this target support VLAs?" Maybe
> > > > > > > > there should be a more specific way to query things about
> > > > > > > > OpenMP devices instead of setting a global flag for the target?
> > > > > > > Actually, the NVPTX and SPIR targets never support VLAs. So I
> > > > > > > felt like it would be more correct to make this a global property
> > > > > > > of the target.
> > > > > > >
> > > > > > > The difference is that the other programming models (OpenCL and
> > > > > > > CUDA) error out immediatelyand regardless of the target because
> > > > > > > this limitation is reflected in the standards that disallow VLAs
> > > > > > > (see SemaType.cpp). For OpenMP we might have target devices that
> > > > > > > support VLA so we shouldn't error out for those.
> > > > > > If you want to make it a global property of the target, that's
> > > > > > fine, but then I don't understand why your diagnostic only fires
> > > > > > when (S.isInOpenMPDeclareTargetContext() ||
> > > > > > S.isInOpenMPTargetExecutionDirective()).
> > > > > That is because of how OpenMP offloading works and how it is
> > > > > implemented in Clang. Consider the following snippet from the added
> > > > > test case:
> > > > > ```lang=c
> > > > > int vla[arg];
> > > > >
> > > > > #pragma omp target map(vla[0:arg])
> > > > > {
> > > > > // more code here...
> > > > > }
> > > > > ```
> > > > >
> > > > > Clang will take the following steps to compile this into a working
> > > > > binary for a GPU:
> > > > > 1. Parse and (semantically) analyze the code as-is for the host and
> > > > > produce LLVM Bitcode.
> > > > > 2. Parse and analyze again the code as-is and generate code for the
> > > > > offloading target, the GPU in this case.
> > > > > 3. Take LLVM Bitcode from 1., generate host binary and embed target
> > > > > binary from 3.
> > > > >
> > > > > `OpenMPIsDevice` will be true for 2., but the complete source code is
> > > > > analyzed. So to not throw errors for the host code, we have to make
> > > > > sure that we are actually generating code for the target device. This
> > > > > is either in a `target` directive or in a `declare target` region.
> > > > > Note that this is quite similar to what CUDA does, only they have
> > > > > `CUDADiagIfDeviceCode` for this logic. If you want me to add
> > > > > something of that kind for OpenMP target devices, I'm fine with that.
> > > > > However for the given case, it's a bit different because this error
> > > > > should only be thrown for target devices that don't support VLAs...
> > > > I see. So the entire translation unit is re-parsed and re-Sema'ed from
> > > > scratch for the target? Which means you need to avoid generating
> > > > errors about things in the outer translation unit that aren't part of
> > > > the target directive that you actually want to compile. I would've
> > > > expected there to be some existing mechanism for that, to be honest, as
> > > > opposed to explicitly trying to suppress target-specific diagnostics
> > > > one by one.
> > > Yes, that is my understanding. For errors, we don't need to take anything
> > > special as the first `cc1` invocation will exit with a non-zero status so
> > > that the driver stops the compilation. For warnings, there seems to be no
> > > mechanism in place as I see them duplicated, even in code that is not
> > > generate for the target device (verified with an unused variable).
> > >
> > > @ABataev @gtbercea Do I miss something here?
> > I'm not aware of any.
> John, target-specific checks require some special flags (like LangOpts.Cuda)
> that are not set when we re-compile the code for OpenMP devices. That's why
> errors are not emitted for the non-target code. But also because of that, we
> need some special OpenMP checks for target-specific code inside the target
> regions. For example, code in lib/Sema/SemaType.cpp, lines 2184, 2185 (see
> this file in this patch) checks for Cuda compilation and prohibits using of
> VLAs in Cuda mode. We also should prohibit using of VLAs in target code for
> NVPTX devices or other devices that do not support VLAs in OpenMP mode.
I think it would be cleaner here, and better for our OpenMP support overall, if
we found a more general way to suppress unwanted diagnostics in the second
invocation for code outside of the target directive. This check (and several
others) would then just implement a more general target feature disabling VLA
support instead of being awkwardly OpenMP-specific.
https://reviews.llvm.org/D39505
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
