Proposed fix: https://reviews.llvm.org/D40275
On Mon, Nov 20, 2017 at 4:13 PM, Artem Belevich <t...@google.com>
wrote:
This change breaks CUDA as clang now reports an error during
device-side compilation when VLA is used in the *host-side* code.
http://lab.llvm.org:8011/builders/clang-cuda-build/builds/15591/steps/ninja%20build%20simple%20CUDA%20tests/logs/stdio
[13]
E.g. I would expect this code to compile successfully, producing
empty device-side binary:
void host_func(int i) {
int vla[i];
}
However it currently fails:
#bin/clang++ --cuda-device-only --cuda-gpu-arch=sm_35 -o vla.o
vla.cu [14]
vla.cu:4:10: error: variable length arrays are not supported for the
current target
int vla[i];
^
1 error generated when compiling for sm_35.
On Sat, Nov 18, 2017 at 1:00 PM, Jonas Hahnfeld via cfe-commits
<cfe-commits@lists.llvm.org> wrote:
Author: hahnfeld
Date: Sat Nov 18 13:00:46 2017
New Revision: 318601
URL: http://llvm.org/viewvc/llvm-project?rev=318601&view=rev [1]
Log:
[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 [2]
Added:
cfe/trunk/test/OpenMP/target_vla_messages.cpp
Modified:
cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
cfe/trunk/include/clang/Basic/TargetInfo.h
cfe/trunk/include/clang/Sema/Sema.h
cfe/trunk/lib/Basic/TargetInfo.cpp
cfe/trunk/lib/Basic/Targets/NVPTX.cpp
cfe/trunk/lib/Basic/Targets/SPIR.h
cfe/trunk/lib/Sema/SemaOpenMP.cpp
cfe/trunk/lib/Sema/SemaType.cpp
Modified: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td?rev=318601&r1=318600&r2=318601&view=diff
[3]
==============================================================================
--- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
(original)
+++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td Sat Nov
18 13:00:46 2017
@@ -141,6 +141,10 @@ def err_vla_decl_has_extern_linkage : Er
"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_ex
"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 {
Modified: cfe/trunk/include/clang/Basic/TargetInfo.h
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/TargetInfo.h?rev=318601&r1=318600&r2=318601&view=diff
[4]
==============================================================================
--- cfe/trunk/include/clang/Basic/TargetInfo.h (original)
+++ cfe/trunk/include/clang/Basic/TargetInfo.h Sat Nov 18 13:00:46
2017
@@ -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() &&
Modified: cfe/trunk/include/clang/Sema/Sema.h
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Sema/Sema.h?rev=318601&r1=318600&r2=318601&view=diff
[5]
==============================================================================
--- cfe/trunk/include/clang/Sema/Sema.h (original)
+++ cfe/trunk/include/clang/Sema/Sema.h Sat Nov 18 13:00:46 2017
@@ -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);
Modified: cfe/trunk/lib/Basic/TargetInfo.cpp
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Basic/TargetInfo.cpp?rev=318601&r1=318600&r2=318601&view=diff
[6]
==============================================================================
--- cfe/trunk/lib/Basic/TargetInfo.cpp (original)
+++ cfe/trunk/lib/Basic/TargetInfo.cpp Sat Nov 18 13:00:46 2017
@@ -31,6 +31,7 @@ TargetInfo::TargetInfo(const llvm::Tripl
// SPARC. These should be overridden by concrete targets as
needed.
BigEndian = !T.isLittleEndian();
TLSSupported = true;
+ VLASupported = true;
NoAsmVariants = false;
HasFloat128 = false;
PointerWidth = PointerAlign = 32;
Modified: cfe/trunk/lib/Basic/Targets/NVPTX.cpp
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Basic/Targets/NVPTX.cpp?rev=318601&r1=318600&r2=318601&view=diff
[7]
==============================================================================
--- cfe/trunk/lib/Basic/Targets/NVPTX.cpp (original)
+++ cfe/trunk/lib/Basic/Targets/NVPTX.cpp Sat Nov 18 13:00:46 2017
@@ -41,6 +41,7 @@ NVPTXTargetInfo::NVPTXTargetInfo(const l
"NVPTX only supports 32- and 64-bit modes.");
TLSSupported = false;
+ VLASupported = false;
AddrSpaceMap = &NVPTXAddrSpaceMap;
UseAddrSpaceMapMangling = true;
Modified: cfe/trunk/lib/Basic/Targets/SPIR.h
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Basic/Targets/SPIR.h?rev=318601&r1=318600&r2=318601&view=diff
[8]
==============================================================================
--- cfe/trunk/lib/Basic/Targets/SPIR.h (original)
+++ cfe/trunk/lib/Basic/Targets/SPIR.h Sat Nov 18 13:00:46 2017
@@ -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;
Modified: cfe/trunk/lib/Sema/SemaOpenMP.cpp
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaOpenMP.cpp?rev=318601&r1=318600&r2=318601&view=diff
[9]
==============================================================================
--- cfe/trunk/lib/Sema/SemaOpenMP.cpp (original)
+++ cfe/trunk/lib/Sema/SemaOpenMP.cpp Sat Nov 18 13:00:46 2017
@@ -1303,6 +1303,17 @@ unsigned Sema::getOpenMPNestingLevel() c
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(Valu
// 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.
Modified: cfe/trunk/lib/Sema/SemaType.cpp
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaType.cpp?rev=318601&r1=318600&r2=318601&view=diff
[10]
==============================================================================
--- cfe/trunk/lib/Sema/SemaType.cpp (original)
+++ cfe/trunk/lib/Sema/SemaType.cpp Sat Nov 18 13:00:46 2017
@@ -2183,6 +2183,12 @@ QualType Sema::BuildArrayType(QualType T
// 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) {
Added: cfe/trunk/test/OpenMP/target_vla_messages.cpp
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/target_vla_messages.cpp?rev=318601&view=auto
[11]
==============================================================================
--- cfe/trunk/test/OpenMP/target_vla_messages.cpp (added)
+++ cfe/trunk/test/OpenMP/target_vla_messages.cpp Sat Nov 18
13:00:46 2017
@@ -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++) ;
+}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits [12]
--
--Artem Belevich
--
--Artem Belevich
Links:
------
[1] http://llvm.org/viewvc/llvm-project?rev=318601&view=rev
[2] https://reviews.llvm.org/D39505
[3]
http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td?rev=318601&r1=318600&r2=318601&view=diff
[4]
http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/TargetInfo.h?rev=318601&r1=318600&r2=318601&view=diff
[5]
http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Sema/Sema.h?rev=318601&r1=318600&r2=318601&view=diff
[6]
http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Basic/TargetInfo.cpp?rev=318601&r1=318600&r2=318601&view=diff
[7]
http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Basic/Targets/NVPTX.cpp?rev=318601&r1=318600&r2=318601&view=diff
[8]
http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Basic/Targets/SPIR.h?rev=318601&r1=318600&r2=318601&view=diff
[9]
http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaOpenMP.cpp?rev=318601&r1=318600&r2=318601&view=diff
[10]
http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaType.cpp?rev=318601&r1=318600&r2=318601&view=diff
[11]
http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/target_vla_messages.cpp?rev=318601&view=auto
[12] http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
[13]
http://lab.llvm.org:8011/builders/clang-cuda-build/builds/15591/steps/ninja%20build%20simple%20CUDA%20tests/logs/stdio
[14] http://vla.cu
