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