r318601 - [OpenMP] Show error if VLAs are not supported
Artem Belevich via cfe-commits
cfe-commits at lists.llvm.org
Mon Nov 20 16:37:30 PST 2017
Proposed fix: https://reviews.llvm.org/D40275
On Mon, Nov 20, 2017 at 4:13 PM, Artem Belevich <tra at 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
>
> 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/Targ
>> etInfo.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/Targ
>> ets/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/Targ
>> ets/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/SemaO
>> penMP.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/SemaT
>> ype.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/ta
>> rget_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
>
--
--Artem Belevich
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20171120/0f787cff/attachment-0001.html>
More information about the cfe-commits
mailing list