[OpenMP] Show error if VLAs are not supported

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.

Differential Revision: https://reviews.llvm.org/D39505

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@318601 91177308-0d34-0410-b5e6-96231b3b80d8

diff --git a/include/clang/Basic/DiagnosticSemaKinds.td b/include/clang/Basic/DiagnosticSemaKinds.td
index 4542fc9194067a37509ed14d2bbd5985ed836d57..def942c1e54f65ccd196cc237ab4a4bcf133d0f1 100644 (file)
--- a/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/include/clang/Basic/DiagnosticSemaKinds.td
@@ -141,6 +141,10 @@ def err_vla_decl_has_extern_linkage : Error<
   "variable length array declaration cannot have 'extern' linkage">;
 def ext_vla_folded_to_constant : Extension<
   "variable length array folded to constant array as an extension">, InGroup<GNUFoldingConstant>;
+def err_vla_unsupported : Error<
+  "variable length arrays are not supported for the current target">;
+def note_vla_unsupported : Note<
+  "variable length arrays are not supported for the current target">;
 
 // C99 variably modified types
 def err_variably_modified_template_arg : Error<
@@ -8985,6 +8989,8 @@ 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_reduction_vla_unsupported : Error<
+  "cannot generate code for reduction on %select{|array section, which requires a }0variable length array">;
 } // end of OpenMP category
 
 let CategoryName = "Related Result Type Issue" in {

diff --git a/include/clang/Basic/TargetInfo.h b/include/clang/Basic/TargetInfo.h
index f80f6b7a6be0a0d0f62ca881c78d9caf983d84c4..da391eedec3708b6027e12a12b86a7e78ecc8bc7 100644 (file)
--- a/include/clang/Basic/TargetInfo.h
+++ b/include/clang/Basic/TargetInfo.h
@@ -60,6 +60,7 @@ protected:
   // 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 @@ public:
     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() &&

diff --git a/include/clang/Sema/Sema.h b/include/clang/Sema/Sema.h
index 8133995c8dfce2885accc99065a1370286db9110..355e20edc036430a2ac69a51b8aa9183542ec499 100644 (file)
--- a/include/clang/Sema/Sema.h
+++ b/include/clang/Sema/Sema.h
@@ -8653,10 +8653,18 @@ public:
                                     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 true if (un)supported features for the current target should be
+  /// diagnosed if OpenMP (offloading) is enabled.
+  bool shouldDiagnoseTargetSupportFromOpenMP() const {
+    return !getLangOpts().OpenMPIsDevice || isInOpenMPDeclareTargetContext() ||
+      isInOpenMPTargetExecutionDirective();
+  }
 
   /// Return the number of captured regions created for an OpenMP directive.
   static int getOpenMPCaptureLevels(OpenMPDirectiveKind Kind);

diff --git a/lib/Basic/TargetInfo.cpp b/lib/Basic/TargetInfo.cpp
index 9ae466b78698d4dd549ee7eb43dd6200c9c68bec..c58ed1e300eb2439f0ed26abb27c00facb55b86f 100644 (file)
--- a/lib/Basic/TargetInfo.cpp
+++ b/lib/Basic/TargetInfo.cpp
@@ -31,6 +31,7 @@ TargetInfo::TargetInfo(const llvm::Triple &T) : TargetOpts(), Triple(T) {
   // SPARC.  These should be overridden by concrete targets as needed.
   BigEndian = !T.isLittleEndian();
   TLSSupported = true;
+  VLASupported = true;
   NoAsmVariants = false;
   HasFloat128 = false;
   PointerWidth = PointerAlign = 32;

diff --git a/lib/Basic/Targets/NVPTX.cpp b/lib/Basic/Targets/NVPTX.cpp
index 3889f0973508ea5e56b24940863619eac75acb13..add3b318aeb6dfaaf30af8fd9548d9f4af9d106a 100644 (file)
--- a/lib/Basic/Targets/NVPTX.cpp
+++ b/lib/Basic/Targets/NVPTX.cpp
@@ -41,6 +41,7 @@ NVPTXTargetInfo::NVPTXTargetInfo(const llvm::Triple &Triple,
          "NVPTX only supports 32- and 64-bit modes.");
 
   TLSSupported = false;
+  VLASupported = false;
   AddrSpaceMap = &NVPTXAddrSpaceMap;
   UseAddrSpaceMapMangling = true;
 

diff --git a/lib/Basic/Targets/SPIR.h b/lib/Basic/Targets/SPIR.h
index e64ad5482908939bdcffc60934e9e47b70410b81..c384d4260ca9e3b19c5aaa43fa6d6ecc0db730e2 100644 (file)
--- a/lib/Basic/Targets/SPIR.h
+++ b/lib/Basic/Targets/SPIR.h
@@ -43,6 +43,7 @@ public:
     assert(getTriple().getEnvironment() == llvm::Triple::UnknownEnvironment &&
            "SPIR target must use unknown environment type");
     TLSSupported = false;
+    VLASupported = false;
     LongWidth = LongAlign = 64;
     AddrSpaceMap = &SPIRAddrSpaceMap;
     UseAddrSpaceMapMangling = true;

diff --git a/lib/Sema/SemaOpenMP.cpp b/lib/Sema/SemaOpenMP.cpp
index 0a48133721dd351d691bef939f548cfc63291fe5..0f349bb961fc76de3d1d82b69ac7a0df98f56f7f 100644 (file)
--- a/lib/Sema/SemaOpenMP.cpp
+++ b/lib/Sema/SemaOpenMP.cpp
@@ -1303,6 +1303,17 @@ unsigned Sema::getOpenMPNestingLevel() const {
   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);
@@ -1315,18 +1326,8 @@ VarDecl *Sema::IsOpenMPCapturedDecl(ValueDecl *D) {
   // 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() ||
@@ -9812,6 +9813,12 @@ static bool ActOnOMPReductionKindClause(
     if ((OASE && !ConstantLengthOASE) ||
         (!OASE && !ASE &&
          D->getType().getNonReferenceType()->isVariablyModifiedType())) {
+      if (!Context.getTargetInfo().isVLASupported() &&
+          S.shouldDiagnoseTargetSupportFromOpenMP()) {
+        S.Diag(ELoc, diag::err_omp_reduction_vla_unsupported) << !!OASE;
+        S.Diag(ELoc, diag::note_vla_unsupported);
+        continue;
+      }
       // For arrays/array sections only:
       // Create pseudo array type for private copy. The size for this array will
       // be generated during codegen.

diff --git a/lib/Sema/SemaType.cpp b/lib/Sema/SemaType.cpp
index ec0d93cf67849a83be5a8f501e32b0cf9c66e773..b00cf7311b977dc244645b3b9e87741f7063d8be 100644 (file)
--- a/lib/Sema/SemaType.cpp
+++ b/lib/Sema/SemaType.cpp
@@ -2183,6 +2183,12 @@ QualType Sema::BuildArrayType(QualType T, ArrayType::ArraySizeModifier ASM,
   // CUDA device code doesn't support VLAs.
   if (getLangOpts().CUDA && T->isVariableArrayType())
     CUDADiagIfDeviceCode(Loc, diag::err_cuda_vla) << CurrentCUDATarget();
+  // Some targets don't support VLAs.
+  if (T->isVariableArrayType() && !Context.getTargetInfo().isVLASupported() &&
+      shouldDiagnoseTargetSupportFromOpenMP()) {
+    Diag(Loc, diag::err_vla_unsupported);
+    return QualType();
+  }
 
   // If this is not C99, extwarn about VLA's and C99 array size modifiers.
   if (!getLangOpts().C99) {

diff --git a/test/OpenMP/target_vla_messages.cpp b/test/OpenMP/target_vla_messages.cpp
new file mode 100644 (file)
index 0000000..c970d0d
--- /dev/null
+++ b/test/OpenMP/target_vla_messages.cpp
@@ -0,0 +1,201 @@
+// 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@+2 {{variable length arrays are not supported for the current target}}
+#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 {{variable length arrays are not supported for the current target}}
+#endif
+#pragma omp parallel reduction(+: a[0:arg])
+  { }
+}
+#pragma omp end declare target
+
+template <typename T>
+void target_template(int arg) {  
+#pragma omp target
+  {
+#ifdef NO_VLA
+    // expected-error@+2 {{variable length arrays are not supported for the current target}}
+#endif
+    T vla[arg];
+  }
+}
+
+void target(int arg) {
+#pragma omp target
+  {
+#ifdef NO_VLA
+    // expected-error@+2 {{variable length arrays are not supported for the current target}}
+#endif
+    int vla[arg];
+  }
+
+#pragma omp target
+  {
+#pragma omp parallel
+    {
+#ifdef NO_VLA
+    // expected-error@+2 {{variable length arrays are not supported for the current target}}
+#endif
+      int vla[arg];
+    }
+  }
+
+  target_template<long>(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 {{variable length arrays are not supported for the current target}}
+#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 {{variable length arrays are not supported for the current target}}
+#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 {{variable length arrays are not supported for the current target}}
+#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 {{variable length arrays are not supported for the current target}}
+#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 {{variable length arrays are not supported for the current target}}
+#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 {{variable length arrays are not supported for the current target}}
+#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 {{variable length arrays are not supported for the current target}}
+#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 {{variable length arrays are not supported for the current target}}
+#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 {{variable length arrays are not supported for the current target}}
+#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++) ;
+}