ABataev added inline comments.
================
Comment at: include/clang/Basic/TargetInfo.h:944
+ /// \brief Whether target supports variable-length arrays.
+ bool isVLASupported() const { return VLASupported; }
+
----------------
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.
https://reviews.llvm.org/D39505
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
