[PATCH] D39505: [OpenMP] Show error if VLAs are not supported

[Jonas Hahnfeld via Phabricator via cfe-commits]

Hahnfeld marked 3 inline comments as done.
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:
> > > 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...


https://reviews.llvm.org/D39505