r318601 - [OpenMP] Show error if VLAs are not supported
Artem Belevich via cfe-commits
cfe-commits at lists.llvm.org
Mon Nov 20 16:13:58 PST 2017
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
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
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 at 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
> 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
>
> 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
> ============================================================
> ==================
> --- 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
> ============================================================
> ==================
> --- 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
> ============================================================
> ==================
> --- 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
> ============================================================
> ==================
> --- 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
> ============================================================
> ==================
> --- 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
> ============================================================
> ==================
> --- 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
> ============================================================
> ==================
> --- 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() && isInOpenMPTargetExecutionDirec
> tive())
> + 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
> ============================================================
> ==================
> --- 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
> ============================================================
> ==================
> --- 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 at +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 at +3 {{cannot generate code for reduction on array
> section, which requires a variable length array}}
> + // expected-note at +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 at +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 at +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 at +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 at +4 {{cannot generate code for reduction on variable
> length array}}
> + // expected-note at +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 at +4 {{cannot generate code for reduction on array
> section, which requires a variable length array}}
> + // expected-note at +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 at +4 {{cannot generate code for reduction on array
> section, which requires a variable length array}}
> + // expected-note at +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 at +4 {{cannot generate code for reduction on variable
> length array}}
> + // expected-note at +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 at +4 {{cannot generate code for reduction on array
> section, which requires a variable length array}}
> + // expected-note at +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 at +4 {{cannot generate code for reduction on array
> section, which requires a variable length array}}
> + // expected-note at +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 at +5 {{cannot generate code for reduction on variable
> length array}}
> + // expected-note at +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 at +5 {{cannot generate code for reduction on array
> section, which requires a variable length array}}
> + // expected-note at +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 at +5 {{cannot generate code for reduction on array
> section, which requires a variable length array}}
> + // expected-note at +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 at lists.llvm.org
> http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
>
--
--Artem Belevich
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20171120/9b34afc4/attachment-0001.html>
More information about the cfe-commits
mailing list