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.bc |
3 | // 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 |
4 | |
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.bc |
6 | // RUN: %clang_cc1 -verify -fopenmp-simd -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 |
7 | |
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.bc |
10 | // 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 |
11 | |
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.bc |
13 | |
14 | #ifndef NO_VLA |
15 | // expected-no-diagnostics |
16 | #endif |
17 | |
18 | #pragma omp declare target |
19 | void declare(int arg) { |
20 | int a[2]; |
21 | #ifdef NO_VLA |
22 | // expected-error@+2 {{variable length arrays are not supported for the current target}} |
23 | #endif |
24 | int vla[arg]; |
25 | } |
26 | |
27 | void 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_VLA |
37 | // 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 | #endif |
40 | #pragma omp parallel reduction(+: a[0:arg]) |
41 | { } |
42 | } |
43 | #pragma omp end declare target |
44 | |
45 | template <typename T> |
46 | void target_template(int arg) { |
47 | #pragma omp target |
48 | { |
49 | #ifdef NO_VLA |
50 | // expected-error@+2 2 {{variable length arrays are not supported for the current target}} |
51 | #endif |
52 | T vla[arg]; |
53 | } |
54 | } |
55 | |
56 | void target(int arg) { |
57 | #pragma omp target |
58 | { |
59 | #ifdef NO_VLA |
60 | // expected-error@+2 {{variable length arrays are not supported for the current target}} |
61 | #endif |
62 | int vla[arg]; |
63 | } |
64 | |
65 | #pragma omp target |
66 | { |
67 | #pragma omp parallel |
68 | { |
69 | #ifdef NO_VLA |
70 | // expected-error@+2 {{variable length arrays are not supported for the current target}} |
71 | #endif |
72 | int vla[arg]; |
73 | } |
74 | } |
75 | |
76 | #ifdef NO_VLA |
77 | // expected-note@+2 {{in instantiation of function template specialization 'target_template<long>' requested here}} |
78 | #endif |
79 | target_template<long>(arg); |
80 | } |
81 | |
82 | void 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_VLA |
91 | // 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 | #endif |
94 | #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_VLA |
107 | // 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 | #endif |
110 | #pragma omp target map(a[0:arg]) |
111 | #pragma omp teams reduction(+: a[0:arg]) |
112 | { } |
113 | |
114 | #ifdef NO_VLA |
115 | // 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 | #endif |
118 | #pragma omp target map(vla[0:arg]) |
119 | #pragma omp teams reduction(+: vla[0:arg]) |
120 | { } |
121 | } |
122 | |
123 | void 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_VLA |
132 | // 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 | #endif |
135 | #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_VLA |
148 | // 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 | #endif |
151 | #pragma omp target map(a[0:arg]) |
152 | #pragma omp parallel reduction(+: a[0:arg]) |
153 | { } |
154 | |
155 | #ifdef NO_VLA |
156 | // 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 | #endif |
159 | #pragma omp target map(vla[0:arg]) |
160 | #pragma omp parallel reduction(+: vla[0:arg]) |
161 | { } |
162 | } |
163 | |
164 | void for_reduction(int arg) { |
165 | int a[2]; |
166 | int vla[arg]; |
167 | |
168 | #pragma omp target map(a) |
169 | #pragma omp parallel |
170 | #pragma omp for reduction(+: a) |
171 | for (int i = 0; i < arg; i++) ; |
172 | |
173 | #ifdef NO_VLA |
174 | // 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 | #endif |
177 | #pragma omp target map(vla) |
178 | #pragma omp parallel |
179 | #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 parallel |
184 | #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 parallel |
189 | #pragma omp for reduction(+: vla[0:2]) |
190 | for (int i = 0; i < arg; i++) ; |
191 | |
192 | #ifdef NO_VLA |
193 | // 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 | #endif |
196 | #pragma omp target map(a[0:arg]) |
197 | #pragma omp parallel |
198 | #pragma omp for reduction(+: a[0:arg]) |
199 | for (int i = 0; i < arg; i++) ; |
200 | |
201 | #ifdef NO_VLA |
202 | // 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 | #endif |
205 | #pragma omp target map(vla[0:arg]) |
206 | #pragma omp parallel |
207 | #pragma omp for reduction(+: vla[0:arg]) |
208 | for (int i = 0; i < arg; i++) ; |
209 | } |
210 | |