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