210 lines · cpp
1// PowerPC supports VLAs.2// 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.bc3// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host-ppc.bc -o %t-ppc-device.ll4 5// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-unknown-unknown -emit-llvm-bc %s -o %t-ppc-host-ppc.bc6// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host-ppc.bc -o %t-ppc-device.ll7 8// Nvidia GPUs don't support VLAs.9// 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.bc10// RUN: %clang_cc1 -verify -DNO_VLA -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host-nvptx.bc -o %t-nvptx-device.ll11 12// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host-nvptx.bc13 14#ifndef NO_VLA15// expected-no-diagnostics16#endif17 18#pragma omp declare target19void declare(int arg) {20 int a[2];21#ifdef NO_VLA22 // expected-error@+2 {{variable length arrays are not supported for the current target}}23#endif24 int vla[arg];25}26 27void declare_parallel_reduction(int arg) {28 int a[2];29 30#pragma omp parallel reduction(+: a)31 { }32 33#pragma omp parallel reduction(+: a[0:2])34 { }35 36#ifdef NO_VLA37 // expected-error@+3 {{cannot generate code for reduction on array section, which requires a variable length array}}38 // expected-note@+2 {{variable length arrays are not supported for the current target}}39#endif40#pragma omp parallel reduction(+: a[0:arg])41 { }42}43#pragma omp end declare target44 45template <typename T>46void target_template(int arg) {47#pragma omp target48 {49#ifdef NO_VLA50 // expected-error@+2 {{variable length arrays are not supported for the current target}}51#endif52 T vla[arg];53 }54}55 56void target(int arg) {57#pragma omp target58 {59#ifdef NO_VLA60 // expected-error@+2 {{variable length arrays are not supported for the current target}}61#endif62 int vla[arg];63 }64 65#pragma omp target66 {67#pragma omp parallel68 {69#ifdef NO_VLA70 // expected-error@+2 {{variable length arrays are not supported for the current target}}71#endif72 int vla[arg];73 }74 }75 76#ifdef NO_VLA77 // expected-note@+2 {{in instantiation of function template specialization 'target_template<long>' requested here}}78#endif79 target_template<long>(arg);80}81 82void teams_reduction(int arg) {83 int a[2];84 int vla[arg];85 86#pragma omp target map(a)87#pragma omp teams reduction(+: a)88 { }89 90#ifdef NO_VLA91 // expected-error@+4 {{cannot generate code for reduction on variable length array}}92 // expected-note@+3 {{variable length arrays are not supported for the current target}}93#endif94#pragma omp target map(vla)95#pragma omp teams reduction(+: vla)96 { }97 98#pragma omp target map(a[0:2])99#pragma omp teams reduction(+: a[0:2])100 { }101 102#pragma omp target map(vla[0:2])103#pragma omp teams reduction(+: vla[0:2])104 { }105 106#ifdef NO_VLA107 // expected-error@+4 {{cannot generate code for reduction on array section, which requires a variable length array}}108 // expected-note@+3 {{variable length arrays are not supported for the current target}}109#endif110#pragma omp target map(a[0:arg])111#pragma omp teams reduction(+: a[0:arg])112 { }113 114#ifdef NO_VLA115 // expected-error@+4 {{cannot generate code for reduction on array section, which requires a variable length array}}116 // expected-note@+3 {{variable length arrays are not supported for the current target}}117#endif118#pragma omp target map(vla[0:arg])119#pragma omp teams reduction(+: vla[0:arg])120 { }121}122 123void parallel_reduction(int arg) {124 int a[2];125 int vla[arg];126 127#pragma omp target map(a)128#pragma omp parallel reduction(+: a)129 { }130 131#ifdef NO_VLA132 // expected-error@+4 {{cannot generate code for reduction on variable length array}}133 // expected-note@+3 {{variable length arrays are not supported for the current target}}134#endif135#pragma omp target map(vla)136#pragma omp parallel reduction(+: vla)137 { }138 139#pragma omp target map(a[0:2])140#pragma omp parallel reduction(+: a[0:2])141 { }142 143#pragma omp target map(vla[0:2])144#pragma omp parallel reduction(+: vla[0:2])145 { }146 147#ifdef NO_VLA148 // expected-error@+4 {{cannot generate code for reduction on array section, which requires a variable length array}}149 // expected-note@+3 {{variable length arrays are not supported for the current target}}150#endif151#pragma omp target map(a[0:arg])152#pragma omp parallel reduction(+: a[0:arg])153 { }154 155#ifdef NO_VLA156 // expected-error@+4 {{cannot generate code for reduction on array section, which requires a variable length array}}157 // expected-note@+3 {{variable length arrays are not supported for the current target}}158#endif159#pragma omp target map(vla[0:arg])160#pragma omp parallel reduction(+: vla[0:arg])161 { }162}163 164void for_reduction(int arg) {165 int a[2];166 int vla[arg];167 168#pragma omp target map(a)169#pragma omp parallel170#pragma omp for reduction(+: a)171 for (int i = 0; i < arg; i++) ;172 173#ifdef NO_VLA174 // expected-error@+5 {{cannot generate code for reduction on variable length array}}175 // expected-note@+4 {{variable length arrays are not supported for the current target}}176#endif177#pragma omp target map(vla)178#pragma omp parallel179#pragma omp for reduction(+: vla)180 for (int i = 0; i < arg; i++) ;181 182#pragma omp target map(a[0:2])183#pragma omp parallel184#pragma omp for reduction(+: a[0:2])185 for (int i = 0; i < arg; i++) ;186 187#pragma omp target map(vla[0:2])188#pragma omp parallel189#pragma omp for reduction(+: vla[0:2])190 for (int i = 0; i < arg; i++) ;191 192#ifdef NO_VLA193 // expected-error@+5 {{cannot generate code for reduction on array section, which requires a variable length array}}194 // expected-note@+4 {{variable length arrays are not supported for the current target}}195#endif196#pragma omp target map(a[0:arg])197#pragma omp parallel198#pragma omp for reduction(+: a[0:arg])199 for (int i = 0; i < arg; i++) ;200 201#ifdef NO_VLA202 // expected-error@+5 {{cannot generate code for reduction on array section, which requires a variable length array}}203 // expected-note@+4 {{variable length arrays are not supported for the current target}}204#endif205#pragma omp target map(vla[0:arg])206#pragma omp parallel207#pragma omp for reduction(+: vla[0:arg])208 for (int i = 0; i < arg; i++) ;209}210