Hahnfeld created this revision.
Herald added a subscriber: jholewinski.
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.
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.c
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++) ;
+}
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: 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/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,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: 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
@@ -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 {
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
