• Home
  • Line#
  • Scopes#
  • Navigate#
  • Raw
  • Download
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
declare(int arg)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 
declare_parallel_reduction(int arg)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>
target_template(int arg)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 
target(int arg)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 
teams_reduction(int arg)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 
parallel_reduction(int arg)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 
for_reduction(int arg)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 #ifdef NO_VLA
210   // expected-error@+3 {{cannot generate code for reduction on array section, which requires a variable length array}}
211   // expected-note@+2 {{variable length arrays are not supported for the current target}}
212 #endif
213 #pragma omp target reduction(+ : vla[0:arg])
214   for (int i = 0; i < arg; i++) ;
215 }
216