[OpenMP] Show error if VLAs are not supported

[Hahnfeld]

Summary

Some target devices (e.g. Nvidia GPUs) don't support dynamic stack
allocation and hence no VLAs. Print errors with description instead
of failing in the backend or generating code that doesn't work.

This patch handles explicit uses of VLAs (local variable in target
or declare target region) or implicitly generated (private) VLAs
for reductions on VLAs or on array sections with non-constant size.

[rjmccall]

include/clang/Basic/TargetInfo.h

944

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?

[Hahnfeld]

include/clang/Basic/TargetInfo.h

944

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.

[rjmccall]

include/clang/Basic/TargetInfo.h

944

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()).

[Hahnfeld]

include/clang/Basic/TargetInfo.h

944

That is because of how OpenMP offloading works and how it is implemented in Clang. Consider the following snippet from the added test case:

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...

[rjmccall]

include/clang/Basic/TargetInfo.h

944

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.

[Hahnfeld]

include/clang/Basic/TargetInfo.h

944

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?

[ABataev]

include/clang/Basic/TargetInfo.h

944

I'm not aware of any.

[ABataev]

include/clang/Basic/TargetInfo.h

944

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.

[rjmccall]

include/clang/Basic/TargetInfo.h

944

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.

[Hahnfeld]

include/clang/Basic/TargetInfo.h

944

I think to get this we would need to make Diag a no-op if (Context.getLangOpts().OpenMPIsDevice && !(isInOpenMPDeclareTargetContext() || isInOpenMPTargetExecutionDirective())). This would ignore all diagnostics outside of the code is really generated in the end...

[rjmccall]

include/clang/Basic/TargetInfo.h

944

I mean, the danger of this approach is that you don't really want to suppress diagnostics for top-level declarations: it can leave you with an invalid AST, and it is not valid to generate IR from an invalid AST.

Sorry for the ignorant questions that follow, but I assume the OpenMP spec must bless this double-translation somehow, and I'd like to understand more about that in order to advise how to proceed. How does OpenMP handle the possibility that the code will be processed substantially differently for different targets? Is there some rule in the spec saying that the code has to expand "the same" in both targets? How does that work when e.g. size_t might have a different size or use a completely different type? More generally, how do expect that this feature will work in the more complicated language modes, like OpenMP + C++?

[Hahnfeld]

include/clang/Basic/TargetInfo.h

944

So if there is an error, the analysis will already fail on the host. I think that guarantees that we don't end up with an invalid AST and will at most suppress duplicate warnings.

Regarding the OpenMP spec: I think the unsatisfying answer is that the spec doesn't say what it expects on that questions. So I think the compiler has to do what seems reasonable...

[rjmccall]

include/clang/Basic/TargetInfo.h

944

Well, what I'm worried about is the possibility that something changes about the translation unit when it's reprocessed for the target — e.g. there's a target-dependent #if that causes an error in the target TU, but not in the original TU, so that the suppressed error is the only reason that the build fails.

If the spec is unclear about this, then we just have to muddle through. Is this "reparse the whole translation unit for the target" the prevailing implementation technique for target directives?

[Hahnfeld]

include/clang/Basic/TargetInfo.h

944

If we are worried about that scenario we have the preserve the current state: Do nothing, diagnose everything and let the user figure out if there is an error in the code.

I can't really comment on what other compilers (GCC, Intel) do, but at least for GCC you compile a complete compiler for the target, so I suppose they kind of do the same...

[ABataev]

include/clang/Basic/TargetInfo.h

944

John, reparsing is required to compile the target-specific code for the particular device. This is an intended behavior. Otherwise we may use some host-specific code on the device. Reparsing allows to reinclude all includes so that the code uses all the definitions for the target instead of those used for the host.
Intel and gcc are doing the same

[rjmccall]

include/clang/Basic/TargetInfo.h

944

Okay, thanks, that's all I was asking.

I am still concerned that this is going to gradually grow to encompass basically every "XX feature is disabled on this target" diagnostic. We could solve this more elegantly by tagging such diagnostics in the .td file and then introducing a diagnostic filter that dropped such diagnostics except when we're inside a target directive.

Even if you don't do that, I'm going to insist that you rewrite the test in SemaType so that it actually fires in normal code if the target disables VLAs. The OpenMP-specific logic should be to suppress the diagnostic if we happen to be in target mode and not within a target directive.

Also, it seems to me that the "are we inside a target directive" check is inadequate to deal with lazy-code-emission features like static functions in C and templates in C++. The diagnostic must be emitted for a violation in a function body if and only if the function is used from within the target directive. That is, unless there's some very coarse language restriction like forbidding external uses in target directives?

Also, shouldn't the "are we in a target directive" be checking that we're in a *matching* target directive? Surely a single OpenMP file can contain target directives for multiple external targets?

[Hahnfeld]

include/clang/Basic/TargetInfo.h

944

Ok, so going forward you suggest something like the following?

// Some targets don't support VLAs.
if (!Context.getTargetInfo().isVLASupported() && T->isVariableArrayType() &&
    (!getLangOpts().OpenMPIsDevice || isInOpenMPDeclareTargetContext() ||
    isInOpenMPTargetExecutionDirective())) {
  Diag(Loc, diag::err_target_no_vla_support);
  return QualType();
}

Doesn't lazy-code emission check the code semantically? Yes, the OpenMP standard (currently) mandates that all external functions that are called must be in a declare target region. And all regions in a file are compiled for all targets, so there shouldn't be a "non-matching" target directive.

[rjmccall]

include/clang/Basic/TargetInfo.h

944

I would check for a VAT before checking whether VLAs are supported, but yes, that's the right idea. Please also consider just adding this check to where we build the VAT in the first place.

We don't (typically) do any extra semantic checks on a function body when it's first used, but if the standard requires all functions that are called to be in a target region, that should be good enough. And ok, I didn't understand that about targets; sounds good.

[Hahnfeld]

Update changes to be generic.

[Hahnfeld]

include/clang/Basic/TargetInfo.h

944

I think I've implemented what we discussed here, please let me know if you disagree. I've also tried to add tests with templates, but declare target is currently broken for that use case, see https://bugs.llvm.org/show_bug.cgi?id=35348.

[rjmccall]

Okay, seems fine. Thanks for putting up with my questions.

include/clang/Basic/DiagnosticSemaKinds.td

def warn_omp_nesting_simd : Warning<
   InGroup<SourceUsesOpenMP>;
 def err_omp_orphaned_device_directive : Error<
   "orphaned 'omp %0' directives are prohibited"
   "; perhaps you forget to enclose the directive into a %select{|||target |teams }1region?">;
 def err_omp_reduction_non_addressable_expression : Error<
   "expected addressable reduction item for the task-based directives">;
 def err_omp_reduction_with_nogroup : Error<
   "'reduction' clause cannot be used with 'nogroup' clause">;
+def err_omp_target_vla : Error<
+  "cannot use variable-length array in a %select{target|declare target}0 region">;
+def err_omp_target_reduction_vla : Error<
+  "cannot generate code for reduction on %select{|array section, which requires a }0variable-length array">;
+def note_omp_target_vla_support : Note<
+  "the target device does not support allocating variable-length arrays">;
 } // end of OpenMP category
 
 let CategoryName = "Related Result Type Issue" in {
 // Objective-C related result type compatibility
 def warn_related_result_type_compatibility_class : Warning<
   "method is expected to return an instance of its class type "
   "%diff{$, but is declared to return $|"
   ", but is declared to return different type}0,1">;

include/clang/Basic/TargetInfo.h

 class TargetInfo : public RefCountedBase<TargetInfo> {
   std::shared_ptr<TargetOptions> TargetOpts;
   llvm::Triple Triple;
 protected:
   // Target values set by the ctor of the actual target implementation.  Default
   // values are specified by the TargetInfo constructor.
   bool BigEndian;
   bool TLSSupported;
+  bool VLASupported;
   bool NoAsmVariants;  // True if {|} are normal characters.
   bool HasFloat128;
   unsigned char PointerWidth, PointerAlign;
   unsigned char BoolWidth, BoolAlign;
   unsigned char IntWidth, IntAlign;
   unsigned char HalfWidth, HalfAlign;
   unsigned char FloatWidth, FloatAlign;
   unsigned char DoubleWidth, DoubleAlign;
public:
   /// \brief Return the maximum alignment (in bits) of a TLS variable
   ///
   /// Gets the maximum alignment (in bits) of a TLS variable on this target.
   /// Returns zero if there is no such constraint.
   unsigned short getMaxTLSAlign() const {
     return MaxTLSAlign;
   }
 
+  /// \brief Whether target supports variable-length arrays.
+  bool isVLASupported() const { return VLASupported; }
+
   /// \brief Whether the target supports SEH __try.
   bool isSEHTrySupported() const {
     return getTriple().isOSWindows() &&
            (getTriple().getArch() == llvm::Triple::x86 ||
             getTriple().getArch() == llvm::Triple::x86_64);
   }
 
   /// \brief Return true if {|} are normal characters in the asm string.

include/clang/Sema/Sema.h

public:
   void ActOnFinishOpenMPDeclareTargetDirective();
   /// Called on correct id-expression from the '#pragma omp declare target'.
   void ActOnOpenMPDeclareTargetName(Scope *CurScope, CXXScopeSpec &ScopeSpec,
                                     const DeclarationNameInfo &Id,
                                     OMPDeclareTargetDeclAttr::MapTypeTy MT,
                                     NamedDeclSetType &SameDirectiveDecls);
   /// Check declaration inside target region.
   void checkDeclIsAllowedInOpenMPTarget(Expr *E, Decl *D);
-  /// Return true inside OpenMP target region.
+  /// Return true inside OpenMP declare target region.
   bool isInOpenMPDeclareTargetContext() const {
     return IsInOpenMPDeclareTargetContext;
   }
+  /// Return true inside OpenMP target region.
+  bool isInOpenMPTargetExecutionDirective() const;
 
   /// Return the number of captured regions created for an OpenMP directive.
   static int getOpenMPCaptureLevels(OpenMPDirectiveKind Kind);
 
   /// \brief Initialization of captured region for OpenMP region.
   void ActOnOpenMPRegionStart(OpenMPDirectiveKind DKind, Scope *CurScope);
   /// \brief End of OpenMP region.
   ///

lib/Basic/TargetInfo.cpp

 static const LangASMap DefaultAddrSpaceMap = {0};
 
 // TargetInfo Constructor.
 TargetInfo::TargetInfo(const llvm::Triple &T) : TargetOpts(), Triple(T) {
   // Set defaults.  Defaults are set for a 32-bit RISC platform, like PPC or
   // SPARC.  These should be overridden by concrete targets as needed.
   BigEndian = !T.isLittleEndian();
   TLSSupported = true;
+  VLASupported = true;
   NoAsmVariants = false;
   HasFloat128 = false;
   PointerWidth = PointerAlign = 32;
   BoolWidth = BoolAlign = 8;
   IntWidth = IntAlign = 32;
   LongWidth = LongAlign = 32;
   LongLongWidth = LongLongAlign = 64;
   SuitableAlign = 64;

lib/Basic/Targets/NVPTX.cpp

 NVPTXTargetInfo::NVPTXTargetInfo(const llvm::Triple &Triple,
                                  const TargetOptions &Opts,
                                  unsigned TargetPointerWidth)
     : TargetInfo(Triple) {
   assert((TargetPointerWidth == 32 || TargetPointerWidth == 64) &&
          "NVPTX only supports 32- and 64-bit modes.");
 
   TLSSupported = false;
+  VLASupported = false;
   AddrSpaceMap = &NVPTXAddrSpaceMap;
   UseAddrSpaceMapMangling = true;
 
   // Define available target features
   // These must be defined in sorted order!
   NoAsmVariants = true;
   GPU = CudaArch::SM_20;
 

lib/Basic/Targets/SPIR.h

 public:
   SPIRTargetInfo(const llvm::Triple &Triple, const TargetOptions &)
       : TargetInfo(Triple) {
     assert(getTriple().getOS() == llvm::Triple::UnknownOS &&
            "SPIR target must use unknown OS");
     assert(getTriple().getEnvironment() == llvm::Triple::UnknownEnvironment &&
            "SPIR target must use unknown environment type");
     TLSSupported = false;
+    VLASupported = false;
     LongWidth = LongAlign = 64;
     AddrSpaceMap = &SPIRAddrSpaceMap;
     UseAddrSpaceMapMangling = true;
     // Define available target features
     // These must be defined in sorted order!
     NoAsmVariants = true;
   }
 

lib/Sema/SemaOpenMP.cpp

bool Sema::IsOpenMPCapturedByRef(ValueDecl *D, unsigned Level) {
   return IsByRef;
 }
 
 unsigned Sema::getOpenMPNestingLevel() const {
   assert(getLangOpts().OpenMP);
   return DSAStack->getNestingLevel();
 }
 
+bool Sema::isInOpenMPTargetExecutionDirective() const {
+  return (isOpenMPTargetExecutionDirective(DSAStack->getCurrentDirective()) &&
+          !DSAStack->isClauseParsingMode()) ||
+         DSAStack->hasDirective(
+             [](OpenMPDirectiveKind K, const DeclarationNameInfo &,
+                SourceLocation) -> bool {
+               return isOpenMPTargetExecutionDirective(K);
+             },
+             false);
+}
+
 VarDecl *Sema::IsOpenMPCapturedDecl(ValueDecl *D) {
   assert(LangOpts.OpenMP && "OpenMP is not allowed");
   D = getCanonicalDecl(D);
 
   // If we are attempting to capture a global variable in a directive with
   // 'target' we return true so that this global is also mapped to the device.
   //
   // FIXME: If the declaration is enclosed in a 'declare target' directive,
   // then it should not be captured. Therefore, an extra check has to be
   // inserted here once support for 'declare target' is added.
   //
   auto *VD = dyn_cast<VarDecl>(D);
-  if (VD && !VD->hasLocalStorage()) {
-    if (isOpenMPTargetExecutionDirective(DSAStack->getCurrentDirective()) &&
-        !DSAStack->isClauseParsingMode())
-      return VD;
-    if (DSAStack->hasDirective(
-            [](OpenMPDirectiveKind K, const DeclarationNameInfo &,
-               SourceLocation) -> bool {
-              return isOpenMPTargetExecutionDirective(K);
-            },
-            false))
-      return VD;
-  }
+  if (VD && !VD->hasLocalStorage() && isInOpenMPTargetExecutionDirective())
 
   if (DSAStack->getCurrentDirective() != OMPD_unknown &&
       (!DSAStack->isClauseParsingMode() ||
        DSAStack->getParentDirective() != OMPD_unknown)) {
     auto &&Info = DSAStack->isLoopControlVariable(D);
     if (Info.first ||
         (VD && VD->hasLocalStorage() &&
          isParallelOrTaskRegion(DSAStack->getCurrentDirective())) ||
if (OASE) {
               PrivateTy, Size, ArrayType::Normal, /*IndexTypeQuals=*/0);
         }
       }
     }
 
     if ((OASE && !ConstantLengthOASE) ||
         (!OASE && !ASE &&
          D->getType().getNonReferenceType()->isVariablyModifiedType())) {
+      if (Context.getLangOpts().OpenMPIsDevice &&
+          !Context.getTargetInfo().isVLASupported() &&
+          (S.isInOpenMPDeclareTargetContext() ||
+           S.isInOpenMPTargetExecutionDirective())) {
+        S.Diag(ELoc, diag::err_omp_target_reduction_vla) << !!OASE;
+        S.Diag(ELoc, diag::note_omp_target_vla_support);
+        continue;
+      }
       // For arrays/array sections only:
       // Create pseudo array type for private copy. The size for this array will
       // be generated during codegen.
       // For array subscripts or single variables Private Ty is the same as Type
       // (type of the variable or single array element).
       PrivateTy = Context.getVariableArrayType(
           Type,
           new (Context) OpaqueValueExpr(ELoc, Context.getSizeType(), VK_RValue),

lib/Sema/SemaType.cpp

QualType Sema::BuildArrayType(QualType T, ArrayType::ArraySizeModifier ASM,
   // OpenCL v1.2 s6.9.d: variable length arrays are not supported.
   if (getLangOpts().OpenCL && T->isVariableArrayType()) {
     Diag(Loc, diag::err_opencl_vla);
     return QualType();
   }
   // CUDA device code doesn't support VLAs.
   if (getLangOpts().CUDA && T->isVariableArrayType())
     CUDADiagIfDeviceCode(Loc, diag::err_cuda_vla) << CurrentCUDATarget();
+  // Some OpenMP target devices don't support VLAs.
+  if (getLangOpts().OpenMPIsDevice &&
+      !Context.getTargetInfo().isVLASupported() && T->isVariableArrayType()) {
+    // Check if we are generating code for the device.
+    bool InDeclareTarget = isInOpenMPDeclareTargetContext();
+    if (InDeclareTarget || isInOpenMPTargetExecutionDirective()) {
+      Diag(Loc, diag::err_omp_target_vla) << InDeclareTarget;
+      Diag(Loc, diag::note_omp_target_vla_support);
+      return QualType();
+    }
+  }
 
   // If this is not C99, extwarn about VLA's and C99 array size modifiers.
   if (!getLangOpts().C99) {
     if (T->isVariableArrayType()) {
       // Prohibit the use of VLAs during template argument deduction.
       if (isSFINAEContext()) {
         Diag(Loc, diag::err_vla_in_sfinae);
         return QualType();

test/OpenMP/target_vla_messages.c

+// PowerPC supports VLAs.
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-unknown-unknown -emit-llvm-bc %s -o %t-ppc-host-ppc.bc
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host-ppc.bc -o %t-ppc-device.ll
+
+// Nvidia GPUs don't support VLAs.
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host-nvptx.bc
+// RUN: %clang_cc1 -verify -DNO_VLA -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host-nvptx.bc -o %t-nvptx-device.ll
+
+#ifndef NO_VLA
+// expected-no-diagnostics
+#endif
+
+#pragma omp declare target
+void declare(int arg) {
+  int a[2];
+#ifdef NO_VLA
+  // expected-error@+3 {{cannot use variable-length array in a declare target region}}
+  // expected-note@+2 {{the target device does not support allocating variable-length arrays}}
+#endif
+  int vla[arg];
+}
+
+void declare_parallel_reduction(int arg) {
+  int a[2];
+
+#pragma omp parallel reduction(+: a)
+  { }
+
+#pragma omp parallel reduction(+: a[0:2])
+  { }
+
+#ifdef NO_VLA
+  // expected-error@+3 {{cannot generate code for reduction on array section, which requires a variable-length array}}
+  // expected-note@+2 {{the target device does not support allocating variable-length arrays}}
+#endif
+#pragma omp parallel reduction(+: a[0:arg])
+  { }
+}
+#pragma omp end declare target
+
+void target(int arg) {
+#pragma omp target
+  {
+#ifdef NO_VLA
+    // expected-error@+3 {{cannot use variable-length array in a target region}}
+    // expected-note@+2 {{the target device does not support allocating variable-length arrays}}
+#endif
+    int vla[arg];
+  }
+
+#pragma omp target
+  {
+#pragma omp parallel
+    {
+#ifdef NO_VLA
+      // expected-error@+3 {{cannot use variable-length array in a target region}}
+      // expected-note@+2 {{the target device does not support allocating variable-length arrays}}
+#endif
+      int vla[arg];
+    }
+  }
+}
+
+void teams_reduction(int arg) {
+  int a[2];
+  int vla[arg];
+
+#pragma omp target map(a)
+#pragma omp teams reduction(+: a)
+  { }
+
+#ifdef NO_VLA
+  // expected-error@+4 {{cannot generate code for reduction on variable-length array}}
+  // expected-note@+3 {{the target device does not support allocating variable-length arrays}}
+#endif
+#pragma omp target map(vla)
+#pragma omp teams reduction(+: vla)
+  { }
+  
+#pragma omp target map(a[0:2])
+#pragma omp teams reduction(+: a[0:2])
+  { }
+
+#pragma omp target map(vla[0:2])
+#pragma omp teams reduction(+: vla[0:2])
+  { }
+
+#ifdef NO_VLA
+  // expected-error@+4 {{cannot generate code for reduction on array section, which requires a variable-length array}}
+  // expected-note@+3 {{the target device does not support allocating variable-length arrays}}
+#endif
+#pragma omp target map(a[0:arg])
+#pragma omp teams reduction(+: a[0:arg])
+  { }
+
+#ifdef NO_VLA
+  // expected-error@+4 {{cannot generate code for reduction on array section, which requires a variable-length array}}
+  // expected-note@+3 {{the target device does not support allocating variable-length arrays}}
+#endif
+#pragma omp target map(vla[0:arg])
+#pragma omp teams reduction(+: vla[0:arg])
+  { }
+}
+
+void parallel_reduction(int arg) {
+  int a[2];
+  int vla[arg];
+
+#pragma omp target map(a)
+#pragma omp parallel reduction(+: a)
+  { }
+
+#ifdef NO_VLA
+  // expected-error@+4 {{cannot generate code for reduction on variable-length array}}
+  // expected-note@+3 {{the target device does not support allocating variable-length arrays}}
+#endif
+#pragma omp target map(vla)
+#pragma omp parallel reduction(+: vla)
+  { }
+
+#pragma omp target map(a[0:2])
+#pragma omp parallel reduction(+: a[0:2])
+  { }
+
+#pragma omp target map(vla[0:2])
+#pragma omp parallel reduction(+: vla[0:2])
+  { }
+
+#ifdef NO_VLA
+  // expected-error@+4 {{cannot generate code for reduction on array section, which requires a variable-length array}}
+  // expected-note@+3 {{the target device does not support allocating variable-length arrays}}
+#endif
+#pragma omp target map(a[0:arg])
+#pragma omp parallel reduction(+: a[0:arg])
+  { }
+
+#ifdef NO_VLA
+  // expected-error@+4 {{cannot generate code for reduction on array section, which requires a variable-length array}}
+  // expected-note@+3 {{the target device does not support allocating variable-length arrays}}
+#endif
+#pragma omp target map(vla[0:arg])
+#pragma omp parallel reduction(+: vla[0:arg])
+  { }
+}
+
+void for_reduction(int arg) {
+  int a[2];
+  int vla[arg];
+
+#pragma omp target map(a)
+#pragma omp parallel
+#pragma omp for reduction(+: a)
+  for (int i = 0; i < arg; i++) ;
+
+#ifdef NO_VLA
+  // expected-error@+5 {{cannot generate code for reduction on variable-length array}}
+  // expected-note@+4 {{the target device does not support allocating variable-length arrays}}
+#endif
+#pragma omp target map(vla)
+#pragma omp parallel
+#pragma omp for reduction(+: vla)
+  for (int i = 0; i < arg; i++) ;
+
+#pragma omp target map(a[0:2])
+#pragma omp parallel
+#pragma omp for reduction(+: a[0:2])
+  for (int i = 0; i < arg; i++) ;
+
+#pragma omp target map(vla[0:2])
+#pragma omp parallel
+#pragma omp for reduction(+: vla[0:2])
+  for (int i = 0; i < arg; i++) ;
+
+#ifdef NO_VLA
+  // expected-error@+5 {{cannot generate code for reduction on array section, which requires a variable-length array}}
+  // expected-note@+4 {{the target device does not support allocating variable-length arrays}}
+#endif
+#pragma omp target map(a[0:arg])
+#pragma omp parallel
+#pragma omp for reduction(+: a[0:arg])
+  for (int i = 0; i < arg; i++) ;
+
+#ifdef NO_VLA
+  // expected-error@+5 {{cannot generate code for reduction on array section, which requires a variable-length array}}
+  // expected-note@+4 {{the target device does not support allocating variable-length arrays}}
+#endif
+#pragma omp target map(vla[0:arg])
+#pragma omp parallel
+#pragma omp for reduction(+: vla[0:arg])
+  for (int i = 0; i < arg; i++) ;
+}