<div dir="ltr">Proposed fix: <a href="https://reviews.llvm.org/D40275">https://reviews.llvm.org/D40275</a></div><div class="gmail_extra"><br><div class="gmail_quote">On Mon, Nov 20, 2017 at 4:13 PM, Artem Belevich <span dir="ltr"><<a href="mailto:tra@google.com" target="_blank">tra@google.com</a>></span> wrote:<br><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex"><div dir="ltr">This change breaks CUDA as clang now reports an error during device-side compilation when VLA is used in the *host-side* code.<div><a href="http://lab.llvm.org:8011/builders/clang-cuda-build/builds/15591/steps/ninja%20build%20simple%20CUDA%20tests/logs/stdio" target="_blank">http://lab.llvm.org:8011/<wbr>builders/clang-cuda-build/<wbr>builds/15591/steps/ninja%<wbr>20build%20simple%20CUDA%<wbr>20tests/logs/stdio</a></div><div><br></div><div>E.g. I would expect this code to compile successfully, producing empty device-side binary:</div><div><br></div><div><div>void host_func(int i) {</div><div> int vla[i];</div><div>}</div></div><div><br></div><div>However it currently fails:</div><div>#bin/clang++ --cuda-device-only --cuda-gpu-arch=sm_35 -o vla.o <a href="http://vla.cu" target="_blank">vla.cu</a><br></div><div><div>vla.cu:4:10: error: variable length arrays are not supported for the current target</div><div> int vla[i];</div><div> ^</div><div>1 error generated when compiling for sm_35.</div></div><div><br></div><div><br><div><br></div></div></div><div class="gmail_extra"><div><div class="h5"><br><div class="gmail_quote">On Sat, Nov 18, 2017 at 1:00 PM, Jonas Hahnfeld via cfe-commits <span dir="ltr"><<a href="mailto:cfe-commits@lists.llvm.org" target="_blank">cfe-commits@lists.llvm.org</a>></span> wrote:<br><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">Author: hahnfeld<br>
Date: Sat Nov 18 13:00:46 2017<br>
New Revision: 318601<br>
<br>
URL: <a href="http://llvm.org/viewvc/llvm-project?rev=318601&view=rev" rel="noreferrer" target="_blank">http://llvm.org/viewvc/llvm-pr<wbr>oject?rev=318601&view=rev</a><br>
Log:<br>
[OpenMP] Show error if VLAs are not supported<br>
<br>
Some target devices (e.g. Nvidia GPUs) don't support dynamic stack<br>
allocation and hence no VLAs. Print errors with description instead<br>
of failing in the backend or generating code that doesn't work.<br>
<br>
This patch handles explicit uses of VLAs (local variable in target<br>
or declare target region) or implicitly generated (private) VLAs<br>
for reductions on VLAs or on array sections with non-constant size.<br>
<br>
Differential Revision: <a href="https://reviews.llvm.org/D39505" rel="noreferrer" target="_blank">https://reviews.llvm.org/D3950<wbr>5</a><br>
<br>
Added:<br>
cfe/trunk/test/OpenMP/target_v<wbr>la_messages.cpp<br>
Modified:<br>
cfe/trunk/include/clang/Basic/<wbr>DiagnosticSemaKinds.td<br>
cfe/trunk/include/clang/Basic/<wbr>TargetInfo.h<br>
cfe/trunk/include/clang/Sema/S<wbr>ema.h<br>
cfe/trunk/lib/Basic/TargetInfo<wbr>.cpp<br>
cfe/trunk/lib/Basic/Targets/NV<wbr>PTX.cpp<br>
cfe/trunk/lib/Basic/Targets/SP<wbr>IR.h<br>
cfe/trunk/lib/Sema/SemaOpenMP.<wbr>cpp<br>
cfe/trunk/lib/Sema/SemaType.cp<wbr>p<br>
<br>
Modified: cfe/trunk/include/clang/Basic/<wbr>DiagnosticSemaKinds.td<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td?rev=318601&r1=318600&r2=318601&view=diff" rel="noreferrer" target="_blank">http://llvm.org/viewvc/llvm-pr<wbr>oject/cfe/trunk/include/clang/<wbr>Basic/DiagnosticSemaKinds.td?<wbr>rev=318601&r1=318600&r2=<wbr>318601&view=diff</a><br>
==============================<wbr>==============================<wbr>==================<br>
--- cfe/trunk/include/clang/Basic/<wbr>DiagnosticSemaKinds.td (original)<br>
+++ cfe/trunk/include/clang/Basic/<wbr>DiagnosticSemaKinds.td Sat Nov 18 13:00:46 2017<br>
@@ -141,6 +141,10 @@ def err_vla_decl_has_extern_linkag<wbr>e : Er<br>
"variable length array declaration cannot have 'extern' linkage">;<br>
def ext_vla_folded_to_constant : Extension<<br>
"variable length array folded to constant array as an extension">, InGroup<GNUFoldingConstant>;<br>
+def err_vla_unsupported : Error<<br>
+ "variable length arrays are not supported for the current target">;<br>
+def note_vla_unsupported : Note<<br>
+ "variable length arrays are not supported for the current target">;<br>
<br>
// C99 variably modified types<br>
def err_variably_modified_template<wbr>_arg : Error<<br>
@@ -8985,6 +8989,8 @@ def err_omp_reduction_non_addressa<wbr>ble_ex<br>
"expected addressable reduction item for the task-based directives">;<br>
def err_omp_reduction_with_nogroup : Error<<br>
"'reduction' clause cannot be used with 'nogroup' clause">;<br>
+def err_omp_reduction_vla_unsuppor<wbr>ted : Error<<br>
+ "cannot generate code for reduction on %select{|array section, which requires a }0variable length array">;<br>
} // end of OpenMP category<br>
<br>
let CategoryName = "Related Result Type Issue" in {<br>
<br>
Modified: cfe/trunk/include/clang/Basic/<wbr>TargetInfo.h<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/TargetInfo.h?rev=318601&r1=318600&r2=318601&view=diff" rel="noreferrer" target="_blank">http://llvm.org/viewvc/llvm-pr<wbr>oject/cfe/trunk/include/clang/<wbr>Basic/TargetInfo.h?rev=318601&<wbr>r1=318600&r2=318601&view=diff</a><br>
==============================<wbr>==============================<wbr>==================<br>
--- cfe/trunk/include/clang/Basic/<wbr>TargetInfo.h (original)<br>
+++ cfe/trunk/include/clang/Basic/<wbr>TargetInfo.h Sat Nov 18 13:00:46 2017<br>
@@ -60,6 +60,7 @@ protected:<br>
// values are specified by the TargetInfo constructor.<br>
bool BigEndian;<br>
bool TLSSupported;<br>
+ bool VLASupported;<br>
bool NoAsmVariants; // True if {|} are normal characters.<br>
bool HasFloat128;<br>
unsigned char PointerWidth, PointerAlign;<br>
@@ -939,6 +940,9 @@ public:<br>
return MaxTLSAlign;<br>
}<br>
<br>
+ /// \brief Whether target supports variable-length arrays.<br>
+ bool isVLASupported() const { return VLASupported; }<br>
+<br>
/// \brief Whether the target supports SEH __try.<br>
bool isSEHTrySupported() const {<br>
return getTriple().isOSWindows() &&<br>
<br>
Modified: cfe/trunk/include/clang/Sema/S<wbr>ema.h<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Sema/Sema.h?rev=318601&r1=318600&r2=318601&view=diff" rel="noreferrer" target="_blank">http://llvm.org/viewvc/llvm-pr<wbr>oject/cfe/trunk/include/clang/<wbr>Sema/Sema.h?rev=318601&r1=<wbr>318600&r2=318601&view=diff</a><br>
==============================<wbr>==============================<wbr>==================<br>
--- cfe/trunk/include/clang/Sema/S<wbr>ema.h (original)<br>
+++ cfe/trunk/include/clang/Sema/S<wbr>ema.h Sat Nov 18 13:00:46 2017<br>
@@ -8653,10 +8653,18 @@ public:<br>
NamedDeclSetType &SameDirectiveDecls);<br>
/// Check declaration inside target region.<br>
void checkDeclIsAllowedInOpenMPTarg<wbr>et(Expr *E, Decl *D);<br>
- /// Return true inside OpenMP target region.<br>
+ /// Return true inside OpenMP declare target region.<br>
bool isInOpenMPDeclareTargetContext<wbr>() const {<br>
return IsInOpenMPDeclareTargetContext<wbr>;<br>
}<br>
+ /// Return true inside OpenMP target region.<br>
+ bool isInOpenMPTargetExecutionDirec<wbr>tive() const;<br>
+ /// Return true if (un)supported features for the current target should be<br>
+ /// diagnosed if OpenMP (offloading) is enabled.<br>
+ bool shouldDiagnoseTargetSupportFro<wbr>mOpenMP() const {<br>
+ return !getLangOpts().OpenMPIsDevice || isInOpenMPDeclareTargetContext<wbr>() ||<br>
+ isInOpenMPTargetExecutionDirec<wbr>tive();<br>
+ }<br>
<br>
/// Return the number of captured regions created for an OpenMP directive.<br>
static int getOpenMPCaptureLevels(OpenMPD<wbr>irectiveKind Kind);<br>
<br>
Modified: cfe/trunk/lib/Basic/TargetInfo<wbr>.cpp<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Basic/TargetInfo.cpp?rev=318601&r1=318600&r2=318601&view=diff" rel="noreferrer" target="_blank">http://llvm.org/viewvc/llvm-pr<wbr>oject/cfe/trunk/lib/Basic/Targ<wbr>etInfo.cpp?rev=318601&r1=31860<wbr>0&r2=318601&view=diff</a><br>
==============================<wbr>==============================<wbr>==================<br>
--- cfe/trunk/lib/Basic/TargetInfo<wbr>.cpp (original)<br>
+++ cfe/trunk/lib/Basic/TargetInfo<wbr>.cpp Sat Nov 18 13:00:46 2017<br>
@@ -31,6 +31,7 @@ TargetInfo::TargetInfo(const llvm::Tripl<br>
// SPARC. These should be overridden by concrete targets as needed.<br>
BigEndian = !T.isLittleEndian();<br>
TLSSupported = true;<br>
+ VLASupported = true;<br>
NoAsmVariants = false;<br>
HasFloat128 = false;<br>
PointerWidth = PointerAlign = 32;<br>
<br>
Modified: cfe/trunk/lib/Basic/Targets/NV<wbr>PTX.cpp<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Basic/Targets/NVPTX.cpp?rev=318601&r1=318600&r2=318601&view=diff" rel="noreferrer" target="_blank">http://llvm.org/viewvc/llvm-pr<wbr>oject/cfe/trunk/lib/Basic/Targ<wbr>ets/NVPTX.cpp?rev=318601&r1=<wbr>318600&r2=318601&view=diff</a><br>
==============================<wbr>==============================<wbr>==================<br>
--- cfe/trunk/lib/Basic/Targets/NV<wbr>PTX.cpp (original)<br>
+++ cfe/trunk/lib/Basic/Targets/NV<wbr>PTX.cpp Sat Nov 18 13:00:46 2017<br>
@@ -41,6 +41,7 @@ NVPTXTargetInfo::NVPTXTargetIn<wbr>fo(const l<br>
"NVPTX only supports 32- and 64-bit modes.");<br>
<br>
TLSSupported = false;<br>
+ VLASupported = false;<br>
AddrSpaceMap = &NVPTXAddrSpaceMap;<br>
UseAddrSpaceMapMangling = true;<br>
<br>
<br>
Modified: cfe/trunk/lib/Basic/Targets/SP<wbr>IR.h<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Basic/Targets/SPIR.h?rev=318601&r1=318600&r2=318601&view=diff" rel="noreferrer" target="_blank">http://llvm.org/viewvc/llvm-pr<wbr>oject/cfe/trunk/lib/Basic/Targ<wbr>ets/SPIR.h?rev=318601&r1=31860<wbr>0&r2=318601&view=diff</a><br>
==============================<wbr>==============================<wbr>==================<br>
--- cfe/trunk/lib/Basic/Targets/SP<wbr>IR.h (original)<br>
+++ cfe/trunk/lib/Basic/Targets/SP<wbr>IR.h Sat Nov 18 13:00:46 2017<br>
@@ -43,6 +43,7 @@ public:<br>
assert(getTriple().getEnviron<wbr>ment() == llvm::Triple::UnknownEnvironme<wbr>nt &&<br>
"SPIR target must use unknown environment type");<br>
TLSSupported = false;<br>
+ VLASupported = false;<br>
LongWidth = LongAlign = 64;<br>
AddrSpaceMap = &SPIRAddrSpaceMap;<br>
UseAddrSpaceMapMangling = true;<br>
<br>
Modified: cfe/trunk/lib/Sema/SemaOpenMP.<wbr>cpp<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaOpenMP.cpp?rev=318601&r1=318600&r2=318601&view=diff" rel="noreferrer" target="_blank">http://llvm.org/viewvc/llvm-pr<wbr>oject/cfe/trunk/lib/Sema/SemaO<wbr>penMP.cpp?rev=318601&r1=318600<wbr>&r2=318601&view=diff</a><br>
==============================<wbr>==============================<wbr>==================<br>
--- cfe/trunk/lib/Sema/SemaOpenMP.<wbr>cpp (original)<br>
+++ cfe/trunk/lib/Sema/SemaOpenMP.<wbr>cpp Sat Nov 18 13:00:46 2017<br>
@@ -1303,6 +1303,17 @@ unsigned Sema::getOpenMPNestingLevel() c<br>
return DSAStack->getNestingLevel();<br>
}<br>
<br>
+bool Sema::isInOpenMPTargetExecutio<wbr>nDirective() const {<br>
+ return (isOpenMPTargetExecutionDirect<wbr>ive(DSAStack->getCurrentDirect<wbr>ive()) &&<br>
+ !DSAStack->isClauseParsingMode<wbr>()) ||<br>
+ DSAStack->hasDirective(<br>
+ [](OpenMPDirectiveKind K, const DeclarationNameInfo &,<br>
+ SourceLocation) -> bool {<br>
+ return isOpenMPTargetExecutionDirecti<wbr>ve(K);<br>
+ },<br>
+ false);<br>
+}<br>
+<br>
VarDecl *Sema::IsOpenMPCapturedDecl(Va<wbr>lueDecl *D) {<br>
assert(LangOpts.OpenMP && "OpenMP is not allowed");<br>
D = getCanonicalDecl(D);<br>
@@ -1315,18 +1326,8 @@ VarDecl *Sema::IsOpenMPCapturedDecl(Va<wbr>lu<br>
// inserted here once support for 'declare target' is added.<br>
//<br>
auto *VD = dyn_cast<VarDecl>(D);<br>
- if (VD && !VD->hasLocalStorage()) {<br>
- if (isOpenMPTargetExecutionDirect<wbr>ive(DSAStack->getCurrentDirect<wbr>ive()) &&<br>
- !DSAStack->isClauseParsingMode<wbr>())<br>
- return VD;<br>
- if (DSAStack->hasDirective(<br>
- [](OpenMPDirectiveKind K, const DeclarationNameInfo &,<br>
- SourceLocation) -> bool {<br>
- return isOpenMPTargetExecutionDirecti<wbr>ve(K);<br>
- },<br>
- false))<br>
- return VD;<br>
- }<br>
+ if (VD && !VD->hasLocalStorage() && isInOpenMPTargetExecutionDirec<wbr>tive())<br>
+ return VD;<br>
<br>
if (DSAStack->getCurrentDirective<wbr>() != OMPD_unknown &&<br>
(!DSAStack->isClauseParsingMo<wbr>de() ||<br>
@@ -9812,6 +9813,12 @@ static bool ActOnOMPReductionKindClause(<br>
if ((OASE && !ConstantLengthOASE) ||<br>
(!OASE && !ASE &&<br>
D->getType().getNonReferenceTy<wbr>pe()->isVariablyModifiedType()<wbr>)) {<br>
+ if (!Context.getTargetInfo().isVL<wbr>ASupported() &&<br>
+ S.shouldDiagnoseTargetSupportF<wbr>romOpenMP()) {<br>
+ S.Diag(ELoc, diag::err_omp_reduction_vla_un<wbr>supported) << !!OASE;<br>
+ S.Diag(ELoc, diag::note_vla_unsupported);<br>
+ continue;<br>
+ }<br>
// For arrays/array sections only:<br>
// Create pseudo array type for private copy. The size for this array will<br>
// be generated during codegen.<br>
<br>
Modified: cfe/trunk/lib/Sema/SemaType.cp<wbr>p<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaType.cpp?rev=318601&r1=318600&r2=318601&view=diff" rel="noreferrer" target="_blank">http://llvm.org/viewvc/llvm-pr<wbr>oject/cfe/trunk/lib/Sema/SemaT<wbr>ype.cpp?rev=318601&r1=318600&<wbr>r2=318601&view=diff</a><br>
==============================<wbr>==============================<wbr>==================<br>
--- cfe/trunk/lib/Sema/SemaType.cp<wbr>p (original)<br>
+++ cfe/trunk/lib/Sema/SemaType.cp<wbr>p Sat Nov 18 13:00:46 2017<br>
@@ -2183,6 +2183,12 @@ QualType Sema::BuildArrayType(QualType T<br>
// CUDA device code doesn't support VLAs.<br>
if (getLangOpts().CUDA && T->isVariableArrayType())<br>
CUDADiagIfDeviceCode(Loc, diag::err_cuda_vla) << CurrentCUDATarget();<br>
+ // Some targets don't support VLAs.<br>
+ if (T->isVariableArrayType() && !Context.getTargetInfo().isVLA<wbr>Supported() &&<br>
+ shouldDiagnoseTargetSupportFro<wbr>mOpenMP()) {<br>
+ Diag(Loc, diag::err_vla_unsupported);<br>
+ return QualType();<br>
+ }<br>
<br>
// If this is not C99, extwarn about VLA's and C99 array size modifiers.<br>
if (!getLangOpts().C99) {<br>
<br>
Added: cfe/trunk/test/OpenMP/target_v<wbr>la_messages.cpp<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/target_vla_messages.cpp?rev=318601&view=auto" rel="noreferrer" target="_blank">http://llvm.org/viewvc/llvm-pr<wbr>oject/cfe/trunk/test/OpenMP/ta<wbr>rget_vla_messages.cpp?rev=3186<wbr>01&view=auto</a><br>
==============================<wbr>==============================<wbr>==================<br>
--- cfe/trunk/test/OpenMP/target_v<wbr>la_messages.cpp (added)<br>
+++ cfe/trunk/test/OpenMP/target_v<wbr>la_messages.cpp Sat Nov 18 13:00:46 2017<br>
@@ -0,0 +1,201 @@<br>
+// PowerPC supports VLAs.<br>
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-u<wbr>nknown-unknown -emit-llvm-bc %s -o %t-ppc-host-ppc.bc<br>
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-u<wbr>nknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host-ppc.bc -o %t-ppc-device.ll<br>
+<br>
+// Nvidia GPUs don't support VLAs.<br>
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidi<wbr>a-cuda -emit-llvm-bc %s -o %t-ppc-host-nvptx.bc<br>
+// RUN: %clang_cc1 -verify -DNO_VLA -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidi<wbr>a-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host-nvptx.bc -o %t-nvptx-device.ll<br>
+<br>
+#ifndef NO_VLA<br>
+// expected-no-diagnostics<br>
+#endif<br>
+<br>
+#pragma omp declare target<br>
+void declare(int arg) {<br>
+ int a[2];<br>
+#ifdef NO_VLA<br>
+ // expected-error@+2 {{variable length arrays are not supported for the current target}}<br>
+#endif<br>
+ int vla[arg];<br>
+}<br>
+<br>
+void declare_parallel_reduction(int arg) {<br>
+ int a[2];<br>
+<br>
+#pragma omp parallel reduction(+: a)<br>
+ { }<br>
+<br>
+#pragma omp parallel reduction(+: a[0:2])<br>
+ { }<br>
+<br>
+#ifdef NO_VLA<br>
+ // expected-error@+3 {{cannot generate code for reduction on array section, which requires a variable length array}}<br>
+ // expected-note@+2 {{variable length arrays are not supported for the current target}}<br>
+#endif<br>
+#pragma omp parallel reduction(+: a[0:arg])<br>
+ { }<br>
+}<br>
+#pragma omp end declare target<br>
+<br>
+template <typename T><br>
+void target_template(int arg) {<br>
+#pragma omp target<br>
+ {<br>
+#ifdef NO_VLA<br>
+ // expected-error@+2 {{variable length arrays are not supported for the current target}}<br>
+#endif<br>
+ T vla[arg];<br>
+ }<br>
+}<br>
+<br>
+void target(int arg) {<br>
+#pragma omp target<br>
+ {<br>
+#ifdef NO_VLA<br>
+ // expected-error@+2 {{variable length arrays are not supported for the current target}}<br>
+#endif<br>
+ int vla[arg];<br>
+ }<br>
+<br>
+#pragma omp target<br>
+ {<br>
+#pragma omp parallel<br>
+ {<br>
+#ifdef NO_VLA<br>
+ // expected-error@+2 {{variable length arrays are not supported for the current target}}<br>
+#endif<br>
+ int vla[arg];<br>
+ }<br>
+ }<br>
+<br>
+ target_template<long>(arg);<br>
+}<br>
+<br>
+void teams_reduction(int arg) {<br>
+ int a[2];<br>
+ int vla[arg];<br>
+<br>
+#pragma omp target map(a)<br>
+#pragma omp teams reduction(+: a)<br>
+ { }<br>
+<br>
+#ifdef NO_VLA<br>
+ // expected-error@+4 {{cannot generate code for reduction on variable length array}}<br>
+ // expected-note@+3 {{variable length arrays are not supported for the current target}}<br>
+#endif<br>
+#pragma omp target map(vla)<br>
+#pragma omp teams reduction(+: vla)<br>
+ { }<br>
+<br>
+#pragma omp target map(a[0:2])<br>
+#pragma omp teams reduction(+: a[0:2])<br>
+ { }<br>
+<br>
+#pragma omp target map(vla[0:2])<br>
+#pragma omp teams reduction(+: vla[0:2])<br>
+ { }<br>
+<br>
+#ifdef NO_VLA<br>
+ // expected-error@+4 {{cannot generate code for reduction on array section, which requires a variable length array}}<br>
+ // expected-note@+3 {{variable length arrays are not supported for the current target}}<br>
+#endif<br>
+#pragma omp target map(a[0:arg])<br>
+#pragma omp teams reduction(+: a[0:arg])<br>
+ { }<br>
+<br>
+#ifdef NO_VLA<br>
+ // expected-error@+4 {{cannot generate code for reduction on array section, which requires a variable length array}}<br>
+ // expected-note@+3 {{variable length arrays are not supported for the current target}}<br>
+#endif<br>
+#pragma omp target map(vla[0:arg])<br>
+#pragma omp teams reduction(+: vla[0:arg])<br>
+ { }<br>
+}<br>
+<br>
+void parallel_reduction(int arg) {<br>
+ int a[2];<br>
+ int vla[arg];<br>
+<br>
+#pragma omp target map(a)<br>
+#pragma omp parallel reduction(+: a)<br>
+ { }<br>
+<br>
+#ifdef NO_VLA<br>
+ // expected-error@+4 {{cannot generate code for reduction on variable length array}}<br>
+ // expected-note@+3 {{variable length arrays are not supported for the current target}}<br>
+#endif<br>
+#pragma omp target map(vla)<br>
+#pragma omp parallel reduction(+: vla)<br>
+ { }<br>
+<br>
+#pragma omp target map(a[0:2])<br>
+#pragma omp parallel reduction(+: a[0:2])<br>
+ { }<br>
+<br>
+#pragma omp target map(vla[0:2])<br>
+#pragma omp parallel reduction(+: vla[0:2])<br>
+ { }<br>
+<br>
+#ifdef NO_VLA<br>
+ // expected-error@+4 {{cannot generate code for reduction on array section, which requires a variable length array}}<br>
+ // expected-note@+3 {{variable length arrays are not supported for the current target}}<br>
+#endif<br>
+#pragma omp target map(a[0:arg])<br>
+#pragma omp parallel reduction(+: a[0:arg])<br>
+ { }<br>
+<br>
+#ifdef NO_VLA<br>
+ // expected-error@+4 {{cannot generate code for reduction on array section, which requires a variable length array}}<br>
+ // expected-note@+3 {{variable length arrays are not supported for the current target}}<br>
+#endif<br>
+#pragma omp target map(vla[0:arg])<br>
+#pragma omp parallel reduction(+: vla[0:arg])<br>
+ { }<br>
+}<br>
+<br>
+void for_reduction(int arg) {<br>
+ int a[2];<br>
+ int vla[arg];<br>
+<br>
+#pragma omp target map(a)<br>
+#pragma omp parallel<br>
+#pragma omp for reduction(+: a)<br>
+ for (int i = 0; i < arg; i++) ;<br>
+<br>
+#ifdef NO_VLA<br>
+ // expected-error@+5 {{cannot generate code for reduction on variable length array}}<br>
+ // expected-note@+4 {{variable length arrays are not supported for the current target}}<br>
+#endif<br>
+#pragma omp target map(vla)<br>
+#pragma omp parallel<br>
+#pragma omp for reduction(+: vla)<br>
+ for (int i = 0; i < arg; i++) ;<br>
+<br>
+#pragma omp target map(a[0:2])<br>
+#pragma omp parallel<br>
+#pragma omp for reduction(+: a[0:2])<br>
+ for (int i = 0; i < arg; i++) ;<br>
+<br>
+#pragma omp target map(vla[0:2])<br>
+#pragma omp parallel<br>
+#pragma omp for reduction(+: vla[0:2])<br>
+ for (int i = 0; i < arg; i++) ;<br>
+<br>
+#ifdef NO_VLA<br>
+ // expected-error@+5 {{cannot generate code for reduction on array section, which requires a variable length array}}<br>
+ // expected-note@+4 {{variable length arrays are not supported for the current target}}<br>
+#endif<br>
+#pragma omp target map(a[0:arg])<br>
+#pragma omp parallel<br>
+#pragma omp for reduction(+: a[0:arg])<br>
+ for (int i = 0; i < arg; i++) ;<br>
+<br>
+#ifdef NO_VLA<br>
+ // expected-error@+5 {{cannot generate code for reduction on array section, which requires a variable length array}}<br>
+ // expected-note@+4 {{variable length arrays are not supported for the current target}}<br>
+#endif<br>
+#pragma omp target map(vla[0:arg])<br>
+#pragma omp parallel<br>
+#pragma omp for reduction(+: vla[0:arg])<br>
+ for (int i = 0; i < arg; i++) ;<br>
+}<br>
<br>
<br>
______________________________<wbr>_________________<br>
cfe-commits mailing list<br>
<a href="mailto:cfe-commits@lists.llvm.org" target="_blank">cfe-commits@lists.llvm.org</a><br>
<a href="http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits" rel="noreferrer" target="_blank">http://lists.llvm.org/cgi-bin/<wbr>mailman/listinfo/cfe-commits</a><br>
</blockquote></div><br><br clear="all"><div><br></div></div></div><span class="HOEnZb"><font color="#888888">-- <br><div class="m_6757778091715542795gmail_signature" data-smartmail="gmail_signature"><div dir="ltr">--Artem Belevich</div></div>
</font></span></div>
</blockquote></div><br><br clear="all"><div><br></div>-- <br><div class="gmail_signature" data-smartmail="gmail_signature"><div dir="ltr">--Artem Belevich</div></div>
</div>