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 > > 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@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@+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 >> > > > > -- > --Artem Belevich > -- --Artem Belevich
_______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits