Index: include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- include/clang/Basic/DiagnosticSemaKinds.td
+++ include/clang/Basic/DiagnosticSemaKinds.td
@@ -8983,6 +8983,12 @@
   "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 {
Index: include/clang/Basic/TargetInfo.h
===================================================================
--- include/clang/Basic/TargetInfo.h
+++ include/clang/Basic/TargetInfo.h
@@ -60,6 +60,7 @@
   // 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;
@@ -939,6 +940,9 @@
     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() &&
Index: include/clang/Sema/Sema.h
===================================================================
--- include/clang/Sema/Sema.h
+++ include/clang/Sema/Sema.h
@@ -8653,10 +8653,12 @@
                                     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);
Index: lib/Basic/TargetInfo.cpp
===================================================================
--- lib/Basic/TargetInfo.cpp
+++ lib/Basic/TargetInfo.cpp
@@ -31,6 +31,7 @@
   // SPARC.  These should be overridden by concrete targets as needed.
   BigEndian = !T.isLittleEndian();
   TLSSupported = true;
+  VLASupported = true;
   NoAsmVariants = false;
   HasFloat128 = false;
   PointerWidth = PointerAlign = 32;
Index: lib/Basic/Targets/NVPTX.cpp
===================================================================
--- lib/Basic/Targets/NVPTX.cpp
+++ lib/Basic/Targets/NVPTX.cpp
@@ -41,6 +41,7 @@
          "NVPTX only supports 32- and 64-bit modes.");
 
   TLSSupported = false;
+  VLASupported = false;
   AddrSpaceMap = &NVPTXAddrSpaceMap;
   UseAddrSpaceMapMangling = true;
 
Index: lib/Basic/Targets/SPIR.h
===================================================================
--- lib/Basic/Targets/SPIR.h
+++ lib/Basic/Targets/SPIR.h
@@ -43,6 +43,7 @@
     assert(getTriple().getEnvironment() == llvm::Triple::UnknownEnvironment &&
            "SPIR target must use unknown environment type");
     TLSSupported = false;
+    VLASupported = false;
     LongWidth = LongAlign = 64;
     AddrSpaceMap = &SPIRAddrSpaceMap;
     UseAddrSpaceMapMangling = true;
Index: lib/Sema/SemaOpenMP.cpp
===================================================================
--- lib/Sema/SemaOpenMP.cpp
+++ lib/Sema/SemaOpenMP.cpp
@@ -1297,6 +1297,17 @@
   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);
@@ -1309,18 +1320,8 @@
   // 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())
+    return VD;
 
   if (DSAStack->getCurrentDirective() != OMPD_unknown &&
       (!DSAStack->isClauseParsingMode() ||
@@ -9713,6 +9714,14 @@
     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.
Index: lib/Sema/SemaType.cpp
===================================================================
--- lib/Sema/SemaType.cpp
+++ lib/Sema/SemaType.cpp
@@ -2183,6 +2183,17 @@
   // 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) {
Index: test/OpenMP/target_vla_messages.c
===================================================================
--- /dev/null
+++ test/OpenMP/target_vla_messages.c
@@ -0,0 +1,191 @@
+// 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++) ;
+}