Index: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td =================================================================== --- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td +++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td @@ -141,6 +141,10 @@ "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; +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 @@ "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 { Index: cfe/trunk/include/clang/Basic/TargetInfo.h =================================================================== --- cfe/trunk/include/clang/Basic/TargetInfo.h +++ cfe/trunk/include/clang/Basic/TargetInfo.h @@ -60,6 +60,7 @@ // 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 @@ 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() && Index: cfe/trunk/include/clang/Sema/Sema.h =================================================================== --- cfe/trunk/include/clang/Sema/Sema.h +++ cfe/trunk/include/clang/Sema/Sema.h @@ -8653,10 +8653,18 @@ 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); Index: cfe/trunk/lib/Basic/TargetInfo.cpp =================================================================== --- cfe/trunk/lib/Basic/TargetInfo.cpp +++ cfe/trunk/lib/Basic/TargetInfo.cpp @@ -31,6 +31,7 @@ // SPARC. These should be overridden by concrete targets as needed. BigEndian = !T.isLittleEndian(); TLSSupported = true; + VLASupported = true; NoAsmVariants = false; HasFloat128 = false; PointerWidth = PointerAlign = 32; Index: cfe/trunk/lib/Basic/Targets/NVPTX.cpp =================================================================== --- cfe/trunk/lib/Basic/Targets/NVPTX.cpp +++ cfe/trunk/lib/Basic/Targets/NVPTX.cpp @@ -41,6 +41,7 @@ "NVPTX only supports 32- and 64-bit modes."); TLSSupported = false; + VLASupported = false; AddrSpaceMap = &NVPTXAddrSpaceMap; UseAddrSpaceMapMangling = true; Index: cfe/trunk/lib/Basic/Targets/SPIR.h =================================================================== --- cfe/trunk/lib/Basic/Targets/SPIR.h +++ cfe/trunk/lib/Basic/Targets/SPIR.h @@ -43,6 +43,7 @@ assert(getTriple().getEnvironment() == llvm::Triple::UnknownEnvironment && "SPIR target must use unknown environment type"); TLSSupported = false; + VLASupported = false; LongWidth = LongAlign = 64; AddrSpaceMap = &SPIRAddrSpaceMap; UseAddrSpaceMapMangling = true; Index: cfe/trunk/lib/Sema/SemaOpenMP.cpp =================================================================== --- cfe/trunk/lib/Sema/SemaOpenMP.cpp +++ cfe/trunk/lib/Sema/SemaOpenMP.cpp @@ -1303,6 +1303,17 @@ 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 @@ // inserted here once support for 'declare target' is added. // auto *VD = dyn_cast(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 @@ 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. Index: cfe/trunk/lib/Sema/SemaType.cpp =================================================================== --- cfe/trunk/lib/Sema/SemaType.cpp +++ cfe/trunk/lib/Sema/SemaType.cpp @@ -2183,6 +2183,12 @@ // 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) { Index: cfe/trunk/test/OpenMP/target_vla_messages.cpp =================================================================== --- cfe/trunk/test/OpenMP/target_vla_messages.cpp +++ cfe/trunk/test/OpenMP/target_vla_messages.cpp @@ -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 +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(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++) ; +}