summaryrefslogtreecommitdiffstats
diff options
context:
space:
mode:
Diffstat
-rw-r--r--clang/include/clang/Basic/DiagnosticSemaKinds.td6
-rw-r--r--clang/include/clang/Basic/TargetInfo.h4
-rw-r--r--clang/include/clang/Sema/Sema.h10
-rw-r--r--clang/lib/Basic/TargetInfo.cpp1
-rw-r--r--clang/lib/Basic/Targets/NVPTX.cpp1
-rw-r--r--clang/lib/Basic/Targets/SPIR.h1
-rw-r--r--clang/lib/Sema/SemaOpenMP.cpp31
-rw-r--r--clang/lib/Sema/SemaType.cpp6
-rw-r--r--clang/test/OpenMP/target_vla_messages.cpp201
9 files changed, 248 insertions, 13 deletions
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 4542fc91940..def942c1e54 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/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/clang/include/clang/Basic/TargetInfo.h b/clang/include/clang/Basic/TargetInfo.h
index f80f6b7a6be..da391eedec3 100644
--- a/clang/include/clang/Basic/TargetInfo.h
+++ b/clang/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/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index 8133995c8df..355e20edc03 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/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/clang/lib/Basic/TargetInfo.cpp b/clang/lib/Basic/TargetInfo.cpp
index 9ae466b7869..c58ed1e300e 100644
--- a/clang/lib/Basic/TargetInfo.cpp
+++ b/clang/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/clang/lib/Basic/Targets/NVPTX.cpp b/clang/lib/Basic/Targets/NVPTX.cpp
index 3889f097350..add3b318aeb 100644
--- a/clang/lib/Basic/Targets/NVPTX.cpp
+++ b/clang/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/clang/lib/Basic/Targets/SPIR.h b/clang/lib/Basic/Targets/SPIR.h
index e64ad548290..c384d4260ca 100644
--- a/clang/lib/Basic/Targets/SPIR.h
+++ b/clang/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/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index 0a48133721d..0f349bb961f 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/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/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp
index ec0d93cf678..b00cf7311b9 100644
--- a/clang/lib/Sema/SemaType.cpp
+++ b/clang/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/clang/test/OpenMP/target_vla_messages.cpp b/clang/test/OpenMP/target_vla_messages.cpp
new file mode 100644
index 00000000000..c970d0def66
--- /dev/null
+++ b/clang/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++) ;
+}
OpenPOWER on IntegriCloud