xref: /llvm-project/clang/test/OpenMP/target_has_device_addr_ast_print.cpp (revision 7c1d9b15eee3a34678addab2bab66f3020ac0753) 
1 // RUN: %clang_cc1 -verify -fopenmp -std=c++11 \
2 // RUN:  -ast-print %s | FileCheck %s
3 
4 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 \
5 // RUN:  -emit-pch -o %t %s
6 
7 // RUN: %clang_cc1 -fopenmp -std=c++11 \
8 // RUN:  -include-pch %t -verify %s -ast-print | FileCheck %s
9 
10 // RUN: %clang_cc1 -verify -fopenmp-simd \
11 // RUN:  -std=c++11 -ast-print %s | FileCheck %s
12 
13 // RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 \
14 // RUN:  -emit-pch -o %t %s
15 
16 // RUN: %clang_cc1 -fopenmp-simd -std=c++11 \
17 // RUN:  -include-pch %t -verify %s -ast-print | FileCheck %s
18 
19 // expected-no-diagnostics
20 
21 #ifndef HEADER
22 #define HEADER
23 
24 struct ST {
25   int *a;
26 };
27 typedef int arr[10];
28 typedef ST STarr[10];
29 struct SA {
30   const int da[5] = { 0 };
31   ST g[10];
32   STarr &rg = g;
33   int i;
34   int &j = i;
35   int *k = &j;
36   int *&z = k;
37   int aa[10];
38   arr &raa = aa;
funcSA39   void func(int arg) {
40 #pragma omp target has_device_addr(k)
41     {}
42 #pragma omp target has_device_addr(z)
43     {}
44 #pragma omp target has_device_addr(aa) // OK
45     {}
46 #pragma omp target has_device_addr(raa) // OK
47     {}
48 #pragma omp target has_device_addr(g) // OK
49     {}
50 #pragma omp target has_device_addr(rg) // OK
51     {}
52 #pragma omp target has_device_addr(da) // OK
53     {}
54   return;
55  }
56 };
57 // CHECK: struct SA
58 // CHECK-NEXT: const int da[5] = {0};
59 // CHECK-NEXT: ST g[10];
60 // CHECK-NEXT: STarr &rg = this->g;
61 // CHECK-NEXT: int i;
62 // CHECK-NEXT: int &j = this->i;
63 // CHECK-NEXT: int *k = &this->j;
64 // CHECK-NEXT: int *&z = this->k;
65 // CHECK-NEXT: int aa[10];
66 // CHECK-NEXT: arr &raa = this->aa;
67 // CHECK-NEXT: func(
68 // CHECK-NEXT: #pragma omp target has_device_addr(this->k)
69 // CHECK-NEXT: {
70 // CHECK-NEXT: }
71 // CHECK-NEXT: #pragma omp target has_device_addr(this->z)
72 // CHECK-NEXT: {
73 // CHECK-NEXT: }
74 // CHECK-NEXT: #pragma omp target has_device_addr(this->aa)
75 // CHECK-NEXT: {
76 // CHECK-NEXT: }
77 // CHECK-NEXT: #pragma omp target has_device_addr(this->raa)
78 // CHECK-NEXT: {
79 // CHECK-NEXT: }
80 // CHECK-NEXT: #pragma omp target has_device_addr(this->g)
81 // CHECK-NEXT: {
82 // CHECK-NEXT: }
83 // CHECK-NEXT: #pragma omp target has_device_addr(this->rg)
84 // CHECK-NEXT: {
85 // CHECK-NEXT: }
86 // CHECK-NEXT: #pragma omp target has_device_addr(this->da)
87 
88 struct SB {
89   unsigned A;
90   unsigned B;
91   float Arr[100];
92   float *Ptr;
fooSB93   float *foo() {
94     return &Arr[0];
95   }
96 };
97 
98 struct SC {
99   unsigned A : 2;
100   unsigned B : 3;
101   unsigned C;
102   unsigned D;
103   float Arr[100];
104   SB S;
105   SB ArrS[100];
106   SB *PtrS;
107   SB *&RPtrS;
108   float *Ptr;
109 
SCSC110   SC(SB *&_RPtrS) : RPtrS(_RPtrS) {}
111 };
112 
113 union SD {
114   unsigned A;
115   float B;
116 };
117 
118 struct S1;
119 extern S1 a;
120 class S2 {
121   mutable int a;
122 public:
S2()123   S2():a(0) { }
S2(S2 & s2)124   S2(S2 &s2):a(s2.a) { }
125   static float S2s;
126   static const float S2sc;
127 };
128 const float S2::S2sc = 0;
129 const S2 b;
130 const S2 ba[5];
131 class S3 {
132   int a;
133 public:
S3()134   S3():a(0) { }
S3(S3 & s3)135   S3(S3 &s3):a(s3.a) { }
136 };
137 const S3 c;
138 const S3 ca[5];
139 extern const int f;
140 class S4 {
141   int a;
142   S4();
143   S4(const S4 &s4);
144 public:
S4(int v)145   S4(int v):a(v) { }
146 };
147 class S5 {
148   int a;
S5()149   S5():a(0) {}
S5(const S5 & s5)150   S5(const S5 &s5):a(s5.a) { }
151 public:
S5(int v)152   S5(int v):a(v) { }
153 };
154 
155 S3 h;
156 #pragma omp threadprivate(h)
157 
158 typedef struct {
159   int a;
160 } S6;
161 
162 template <typename T>
tmain(T argc)163 T tmain(T argc) {
164   const T da[5] = { 0 };
165   S6 h[10];
166   auto &rh = h;
167   T i;
168   T &j = i;
169   T *k = &j;
170   T *&z = k;
171   T aa[10];
172   auto &raa = aa;
173 #pragma omp target has_device_addr(k)
174   {}
175 #pragma omp target has_device_addr(z)
176   {}
177 #pragma omp target has_device_addr(aa)
178   {}
179 #pragma omp target has_device_addr(raa)
180   {}
181 #pragma omp target has_device_addr(h)
182   {}
183 #pragma omp target has_device_addr(rh)
184   {}
185 #pragma omp target has_device_addr(da)
186   {}
187   return 0;
188 }
189 
190 // CHECK: template<> int tmain<int>(int argc) {
191 // CHECK-NEXT: const int da[5] = {0};
192 // CHECK-NEXT: S6 h[10];
193 // CHECK-NEXT: auto &rh = h;
194 // CHECK-NEXT: int i;
195 // CHECK-NEXT: int &j = i;
196 // CHECK-NEXT: int *k = &j;
197 // CHECK-NEXT: int *&z = k;
198 // CHECK-NEXT: int aa[10];
199 // CHECK-NEXT: auto &raa = aa;
200 // CHECK-NEXT: #pragma omp target has_device_addr(k)
201 // CHECK-NEXT: {
202 // CHECK-NEXT: }
203 // CHECK-NEXT: #pragma omp target has_device_addr(z)
204 // CHECK-NEXT: {
205 // CHECK-NEXT: }
206 // CHECK-NEXT: #pragma omp target has_device_addr(aa)
207 // CHECK-NEXT: {
208 // CHECK-NEXT: }
209 // CHECK-NEXT: #pragma omp target has_device_addr(raa)
210 // CHECK-NEXT: {
211 // CHECK-NEXT: }
212 // CHECK-NEXT: #pragma omp target has_device_addr(h)
213 // CHECK-NEXT: {
214 // CHECK-NEXT: }
215 // CHECK-NEXT: #pragma omp target has_device_addr(rh)
216 // CHECK-NEXT: {
217 // CHECK-NEXT: }
218 // CHECK-NEXT: #pragma omp target has_device_addr(da)
219 
220 // CHECK: template<> int *tmain<int *>(int *argc) {
221 // CHECK-NEXT: int *const da[5] = {0};
222 // CHECK-NEXT: S6 h[10];
223 // CHECK-NEXT: auto &rh = h;
224 // CHECK-NEXT: int *i;
225 // CHECK-NEXT: int *&j = i;
226 // CHECK-NEXT: int **k = &j;
227 // CHECK-NEXT: int **&z = k;
228 // CHECK-NEXT: int *aa[10];
229 // CHECK-NEXT: auto &raa = aa;
230 // CHECK-NEXT: #pragma omp target has_device_addr(k)
231 // CHECK-NEXT: {
232 // CHECK-NEXT: }
233 // CHECK-NEXT: #pragma omp target has_device_addr(z)
234 // CHECK-NEXT: {
235 // CHECK-NEXT: }
236 // CHECK-NEXT: #pragma omp target has_device_addr(aa)
237 // CHECK-NEXT: {
238 // CHECK-NEXT: }
239 // CHECK-NEXT: #pragma omp target has_device_addr(raa)
240 // CHECK-NEXT: {
241 // CHECK-NEXT: }
242 // CHECK-NEXT: #pragma omp target has_device_addr(h)
243 // CHECK-NEXT: {
244 // CHECK-NEXT: }
245 // CHECK-NEXT: #pragma omp target has_device_addr(rh)
246 // CHECK-NEXT: {
247 // CHECK-NEXT: }
248 // CHECK-NEXT: #pragma omp target has_device_addr(da)
249 
250 // CHECK-LABEL: int main(int argc, char **argv) {
main(int argc,char ** argv)251 int main(int argc, char **argv) {
252   const int da[5] = { 0 };
253   S6 h[10];
254   auto &rh = h;
255   int i;
256   int &j = i;
257   int *k = &j;
258   int *&z = k;
259   int aa[10];
260   auto &raa = aa;
261 // CHECK-NEXT: const int da[5] = {0};
262 // CHECK-NEXT: S6 h[10];
263 // CHECK-NEXT: auto &rh = h;
264 // CHECK-NEXT: int i;
265 // CHECK-NEXT: int &j = i;
266 // CHECK-NEXT: int *k = &j;
267 // CHECK-NEXT: int *&z = k;
268 // CHECK-NEXT: int aa[10];
269 // CHECK-NEXT: auto &raa = aa;
270 #pragma omp target has_device_addr(k)
271 // CHECK-NEXT: #pragma omp target has_device_addr(k)
272   {}
273 // CHECK-NEXT: {
274 // CHECK-NEXT: }
275 #pragma omp target has_device_addr(z)
276 // CHECK-NEXT: #pragma omp target has_device_addr(z)
277   {}
278 // CHECK-NEXT: {
279 // CHECK-NEXT: }
280 #pragma omp target has_device_addr(aa)
281 // CHECK-NEXT: #pragma omp target has_device_addr(aa)
282   {}
283 // CHECK-NEXT: {
284 // CHECK-NEXT: }
285 #pragma omp target has_device_addr(raa)
286 // CHECK-NEXT: #pragma omp target has_device_addr(raa)
287   {}
288 // CHECK-NEXT: {
289 // CHECK-NEXT: }
290 #pragma omp target has_device_addr(h)
291 // CHECK-NEXT: #pragma omp target has_device_addr(h)
292   {}
293 // CHECK-NEXT: {
294 // CHECK-NEXT: }
295 #pragma omp target has_device_addr(rh)
296 // CHECK-NEXT: #pragma omp target has_device_addr(rh)
297   {}
298 // CHECK-NEXT: {
299 // CHECK-NEXT: }
300 #pragma omp target has_device_addr(da)
301 // CHECK-NEXT: #pragma omp target has_device_addr(da)
302   {}
303 // CHECK-NEXT: {
304 // CHECK-NEXT: }
305 // CHECK-NEXT: #pragma omp target has_device_addr(da[1:3])
306 // CHECK-NEXT: {
307 // CHECK-NEXT: }
308 #pragma omp target has_device_addr(da[1:3])
309   {}
310   return tmain<int>(argc) + *tmain<int *>(&argc);
311 }
312 
313 struct SomeKernel {
314   int targetDev;
315   float devPtr;
316   SomeKernel();
317   ~SomeKernel();
318 
319   template<unsigned int nRHS>
applySomeKernel320   void apply() {
321     #pragma omp target has_device_addr(devPtr) device(targetDev)
322     {
323     }
324 // CHECK:  #pragma omp target has_device_addr(this->devPtr) device(this->targetDev)
325 // CHECK-NEXT: {
326 // CHECK-NEXT: }
327   }
328 // CHECK: template<> void apply<32U>() {
329 // CHECK: #pragma omp target has_device_addr(this->devPtr) device(this->targetDev)
330 // CHECK-NEXT: {
331 // CHECK-NEXT: }
332 };
333 
use_template()334 void use_template() {
335   SomeKernel aKern;
336   aKern.apply<32>();
337 }
338 #endif
339