Hahnfeld updated this revision to Diff 123364.
Hahnfeld added a comment.
Update changes to be generic.
https://reviews.llvm.org/D39505
Files:
include/clang/Basic/DiagnosticSemaKinds.td
include/clang/Basic/TargetInfo.h
include/clang/Sema/Sema.h
lib/Basic/TargetInfo.cpp
lib/Basic/Targets/NVPTX.cpp
lib/Basic/Targets/SPIR.h
lib/Sema/SemaOpenMP.cpp
lib/Sema/SemaType.cpp
test/OpenMP/target_vla_messages.cpp
Index: test/OpenMP/target_vla_messages.cpp
===================================================================
--- /dev/null
+++ 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++) ;
+}
Index: lib/Sema/SemaType.cpp
===================================================================
--- lib/Sema/SemaType.cpp
+++ lib/Sema/SemaType.cpp
@@ -2183,6 +2183,12 @@
// 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) {
Index: lib/Sema/SemaOpenMP.cpp
===================================================================
--- lib/Sema/SemaOpenMP.cpp
+++ lib/Sema/SemaOpenMP.cpp
@@ -1303,6 +1303,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);
@@ -1315,18 +1326,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() ||
@@ -9801,6 +9802,12 @@
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.
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/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/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: include/clang/Sema/Sema.h
===================================================================
--- include/clang/Sema/Sema.h
+++ include/clang/Sema/Sema.h
@@ -8653,10 +8653,18 @@
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);
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/Basic/DiagnosticSemaKinds.td
===================================================================
--- include/clang/Basic/DiagnosticSemaKinds.td
+++ include/clang/Basic/DiagnosticSemaKinds.td
@@ -141,6 +141,10 @@
"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<
@@ -8983,6 +8987,8 @@
"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 {
_______________________________________________
cfe-commits mailing list
[email protected]
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
