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 {{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   target_template<long>(arg);
77 }
78 
79 void teams_reduction(int arg) {
80   int a[2];
81   int vla[arg];
82 
83 #pragma omp target map(a)
84 #pragma omp teams reduction(+: a)
85   { }
86 
87 #ifdef NO_VLA
88   // expected-error@+4 {{cannot generate code for reduction on variable length array}}
89   // expected-note@+3 {{variable length arrays are not supported for the current target}}
90 #endif
91 #pragma omp target map(vla)
92 #pragma omp teams reduction(+: vla)
93   { }
94 
95 #pragma omp target map(a[0:2])
96 #pragma omp teams reduction(+: a[0:2])
97   { }
98 
99 #pragma omp target map(vla[0:2])
100 #pragma omp teams reduction(+: vla[0:2])
101   { }
102 
103 #ifdef NO_VLA
104   // expected-error@+4 {{cannot generate code for reduction on array section, which requires a variable length array}}
105   // expected-note@+3 {{variable length arrays are not supported for the current target}}
106 #endif
107 #pragma omp target map(a[0:arg])
108 #pragma omp teams reduction(+: a[0:arg])
109   { }
110 
111 #ifdef NO_VLA
112   // expected-error@+4 {{cannot generate code for reduction on array section, which requires a variable length array}}
113   // expected-note@+3 {{variable length arrays are not supported for the current target}}
114 #endif
115 #pragma omp target map(vla[0:arg])
116 #pragma omp teams reduction(+: vla[0:arg])
117   { }
118 }
119 
120 void parallel_reduction(int arg) {
121   int a[2];
122   int vla[arg];
123 
124 #pragma omp target map(a)
125 #pragma omp parallel reduction(+: a)
126   { }
127 
128 #ifdef NO_VLA
129   // expected-error@+4 {{cannot generate code for reduction on variable length array}}
130   // expected-note@+3 {{variable length arrays are not supported for the current target}}
131 #endif
132 #pragma omp target map(vla)
133 #pragma omp parallel reduction(+: vla)
134   { }
135 
136 #pragma omp target map(a[0:2])
137 #pragma omp parallel reduction(+: a[0:2])
138   { }
139 
140 #pragma omp target map(vla[0:2])
141 #pragma omp parallel reduction(+: vla[0:2])
142   { }
143 
144 #ifdef NO_VLA
145   // expected-error@+4 {{cannot generate code for reduction on array section, which requires a variable length array}}
146   // expected-note@+3 {{variable length arrays are not supported for the current target}}
147 #endif
148 #pragma omp target map(a[0:arg])
149 #pragma omp parallel reduction(+: a[0:arg])
150   { }
151 
152 #ifdef NO_VLA
153   // expected-error@+4 {{cannot generate code for reduction on array section, which requires a variable length array}}
154   // expected-note@+3 {{variable length arrays are not supported for the current target}}
155 #endif
156 #pragma omp target map(vla[0:arg])
157 #pragma omp parallel reduction(+: vla[0:arg])
158   { }
159 }
160 
161 void for_reduction(int arg) {
162   int a[2];
163   int vla[arg];
164 
165 #pragma omp target map(a)
166 #pragma omp parallel
167 #pragma omp for reduction(+: a)
168   for (int i = 0; i < arg; i++) ;
169 
170 #ifdef NO_VLA
171   // expected-error@+5 {{cannot generate code for reduction on variable length array}}
172   // expected-note@+4 {{variable length arrays are not supported for the current target}}
173 #endif
174 #pragma omp target map(vla)
175 #pragma omp parallel
176 #pragma omp for reduction(+: vla)
177   for (int i = 0; i < arg; i++) ;
178 
179 #pragma omp target map(a[0:2])
180 #pragma omp parallel
181 #pragma omp for reduction(+: a[0:2])
182   for (int i = 0; i < arg; i++) ;
183 
184 #pragma omp target map(vla[0:2])
185 #pragma omp parallel
186 #pragma omp for reduction(+: vla[0:2])
187   for (int i = 0; i < arg; i++) ;
188 
189 #ifdef NO_VLA
190   // expected-error@+5 {{cannot generate code for reduction on array section, which requires a variable length array}}
191   // expected-note@+4 {{variable length arrays are not supported for the current target}}
192 #endif
193 #pragma omp target map(a[0:arg])
194 #pragma omp parallel
195 #pragma omp for reduction(+: a[0:arg])
196   for (int i = 0; i < arg; i++) ;
197 
198 #ifdef NO_VLA
199   // expected-error@+5 {{cannot generate code for reduction on array section, which requires a variable length array}}
200   // expected-note@+4 {{variable length arrays are not supported for the current target}}
201 #endif
202 #pragma omp target map(vla[0:arg])
203 #pragma omp parallel
204 #pragma omp for reduction(+: vla[0:arg])
205   for (int i = 0; i < arg; i++) ;
206 }
207