xref: /llvm-project/clang/test/OpenMP/target_codegen_registration.cpp (revision 13dcc95dcd4999ff99f2de89d881f1aed5b21709) 
1 // Test host codegen.
2 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s
3 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
4 // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
5 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s
6 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
7 // RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
8 
9 // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
10 // RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
11 // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
12 // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
13 // RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
14 // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
15 // SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
16 
17 // Test target codegen - host bc file has to be created first.
18 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc
19 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s -check-prefix=TCHECK
20 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s
21 // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s -check-prefix=TCHECK
22 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc
23 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s -check-prefix=TCHECK
24 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s
25 // RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s -check-prefix=TCHECK
26 
27 // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc
28 // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck --check-prefix SIMD-ONLY1 %s
29 // RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s
30 // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY1 %s
31 // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc
32 // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck --check-prefix SIMD-ONLY1 %s
33 // RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s
34 // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY1 %s
35 // SIMD-ONLY1-NOT: {{__kmpc|__tgt}}
36 
37 // Check that no target code is emitted if no omptests flag was provided.
38 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s -check-prefix=CHECK-NTARGET
39 
40 // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY2 %s
41 // SIMD-ONLY2-NOT: {{__kmpc|__tgt}}
42 
43 // expected-no-diagnostics
44 #ifndef HEADER
45 #define HEADER
46 
47 // CHECK-DAG: [[SA:%.+]] = type { [4 x i32] }
48 // CHECK-DAG: [[SB:%.+]] = type { [8 x i32] }
49 // CHECK-DAG: [[SC:%.+]] = type { [16 x i32] }
50 // CHECK-DAG: [[SD:%.+]] = type { [32 x i32] }
51 // CHECK-DAG: [[SE:%.+]] = type { [64 x i32] }
52 // CHECK-DAG: [[ST1:%.+]] = type { [228 x i32] }
53 // CHECK-DAG: [[ST2:%.+]] = type { [1128 x i32] }
54 // CHECK-DAG: [[ENTTY:%.+]] = type { i64, i16, i16, i32, ptr, ptr, i64, i64, ptr }
55 
56 // TCHECK: [[ENTTY:%.+]] = type { i64, i16, i16, i32, ptr, ptr, i64, i64, ptr }
57 
58 // CHECK-DAG: [[A1:@.+]] = internal global [[SA]]
59 // CHECK-DAG: [[A2:@.+]] ={{.*}} global [[SA]]
60 // CHECK-DAG: [[B1:@.+]] ={{.*}} global [[SB]]
61 // CHECK-DAG: [[B2:@.+]] ={{.*}} global [[SB]]
62 // CHECK-DAG: [[C1:@.+]] = internal global [[SC]]
63 // CHECK-DAG: [[D1:@.+]] ={{.*}} global [[SD]]
64 // CHECK-DAG: [[E1:@.+]] ={{.*}} global [[SE]]
65 // CHECK-DAG: [[T1:@.+]] ={{.*}} global [[ST1]]
66 // CHECK-DAG: [[T2:@.+]] ={{.*}} global [[ST2]]
67 
68 // CHECK-NTARGET-DAG: [[SA:%.+]] = type { [4 x i32] }
69 // CHECK-NTARGET-DAG: [[SB:%.+]] = type { [8 x i32] }
70 // CHECK-NTARGET-DAG: [[SC:%.+]] = type { [16 x i32] }
71 // CHECK-NTARGET-DAG: [[SD:%.+]] = type { [32 x i32] }
72 // CHECK-NTARGET-DAG: [[SE:%.+]] = type { [64 x i32] }
73 // CHECK-NTARGET-DAG: [[ST1:%.+]] = type { [228 x i32] }
74 // CHECK-NTARGET-DAG: [[ST2:%.+]] = type { [1128 x i32] }
75 // CHECK-NTARGET-NOT: type { ptr,
76 // CHECK-NTARGET-NOT: type { i32,
77 
78 // We have 7 target regions
79 
80 // CHECK-DAG: {{@.+}} = weak{{.*}} constant i8 0
81 // TCHECK-NOT: {{@.+}} = weak{{.*}} constant i8 0
82 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 4]
83 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 800]
84 // CHECK-DAG: {{@.+}} = weak{{.*}} constant i8 0
85 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 4]
86 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 800]
87 // CHECK-DAG: {{@.+}} = weak{{.*}} constant i8 0
88 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 4]
89 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 800]
90 // CHECK-DAG: {{@.+}} = weak{{.*}} constant i8 0
91 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 4]
92 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 800]
93 // CHECK-DAG: {{@.+}} = weak{{.*}} constant i8 0
94 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 4]
95 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 800]
96 // CHECK-DAG: {{@.+}} = weak{{.*}} constant i8 0
97 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 4]
98 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 800]
99 // CHECK-DAG: {{@.+}} = weak{{.*}} constant i8 0
100 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 4]
101 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 800]
102 // CHECK-DAG: {{@.+}} = weak{{.*}} constant i8 0
103 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 4]
104 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 800]
105 // CHECK-DAG: {{@.+}} = weak{{.*}} constant i8 0
106 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 4]
107 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 800]
108 // CHECK-DAG: {{@.+}} = weak{{.*}} constant i8 0
109 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 4]
110 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 800]
111 // CHECK-DAG: {{@.+}} = weak{{.*}} constant i8 0
112 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 4]
113 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 800]
114 // CHECK-DAG: {{@.+}} = weak{{.*}} constant i8 0
115 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 4]
116 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i64] [i64 800]
117 
118 // CHECK-NTARGET-NOT: weak constant i8 0
119 // CHECK-NTARGET-NOT: private unnamed_addr constant [1 x i
120 
121 // CHECK-DAG: [[NAMEPTR1:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME1:__omp_offloading_[0-9a-f]+_[0-9a-f]+__Z.+_l[0-9]+]]\00"
122 // CHECK-DAG: [[NAMEPTR2:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME2:.+]]\00"
123 // CHECK-DAG: [[NAMEPTR3:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME3:.+]]\00"
124 // CHECK-DAG: [[NAMEPTR4:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME4:.+]]\00"
125 // CHECK-DAG: [[NAMEPTR5:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME5:.+]]\00"
126 // CHECK-DAG: [[NAMEPTR6:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME6:.+]]\00"
127 // CHECK-DAG: [[NAMEPTR7:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME7:.+]]\00"
128 // CHECK-DAG: [[NAMEPTR8:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME8:.+]]\00"
129 // CHECK-DAG: [[NAMEPTR9:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME9:.+]]\00"
130 // CHECK-DAG: [[NAMEPTR10:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME10:.+]]\00"
131 // CHECK-DAG: [[NAMEPTR11:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME11:.+]]\00"
132 // CHECK-DAG: [[NAMEPTR12:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME12:.+]]\00"
133 
134 // TCHECK-DAG: [[NAMEPTR1:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME1:__omp_offloading_[0-9a-f]+_[0-9a-f]+__Z.+_l[0-9]+]]\00"
135 // TCHECK-DAG: [[NAMEPTR2:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME2:.+]]\00"
136 // TCHECK-DAG: [[NAMEPTR3:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME3:.+]]\00"
137 // TCHECK-DAG: [[NAMEPTR4:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME4:.+]]\00"
138 // TCHECK-DAG: [[NAMEPTR5:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME5:.+]]\00"
139 // TCHECK-DAG: [[NAMEPTR6:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME6:.+]]\00"
140 // TCHECK-DAG: [[NAMEPTR7:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME7:.+]]\00"
141 // TCHECK-DAG: [[NAMEPTR8:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME8:.+]]\00"
142 // TCHECK-DAG: [[NAMEPTR9:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME9:.+]]\00"
143 // TCHECK-DAG: [[NAMEPTR10:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME10:.+]]\00"
144 // TCHECK-DAG: [[NAMEPTR11:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME11:.+]]\00"
145 // TCHECK-DAG: [[NAMEPTR12:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME12:.+]]\00"
146 
147 // We have 4 initializers, one for the 500 priority, another one for 501, or more for the default priority, and the last one for the offloading registration function.
148 // CHECK: @llvm.global_ctors = appending global [3 x { i32, ptr, ptr }] [
149 // CHECK-SAME: { i32, ptr, ptr } { i32 500, ptr [[P500:@[^,]+]], ptr null },
150 // CHECK-SAME: { i32, ptr, ptr } { i32 501, ptr [[P501:@[^,]+]], ptr null },
151 // CHECK-SAME: { i32, ptr, ptr } { i32 65535, ptr [[PMAX:@[^,]+]], ptr null }
152 
153 // CHECK-NTARGET: @llvm.global_ctors = appending global [3   x { i32, ptr, ptr }] [
154 
155 extern int *R;
156 
157 struct SA {
158   int arr[4];
159   void foo() {
160     int a = *R;
161     a += 1;
162     *R = a;
163   }
164   SA() {
165     int a = *R;
166     a += 2;
167     *R = a;
168   }
169   ~SA() {
170     int a = *R;
171     a += 3;
172     *R = a;
173   }
174 };
175 
176 struct SB {
177   int arr[8];
178   void foo() {
179     int a = *R;
180     #pragma omp target
181     a += 4;
182     *R = a;
183   }
184   SB() {
185     int a = *R;
186     a += 5;
187     *R = a;
188   }
189   ~SB() {
190     int a = *R;
191     a += 6;
192     *R = a;
193   }
194 };
195 
196 struct SC {
197   int arr[16];
198   void foo() {
199     int a = *R;
200     a += 7;
201     *R = a;
202   }
203   SC() {
204     int a = *R;
205     #pragma omp target
206     a += 8;
207     *R = a;
208   }
209   ~SC() {
210     int a = *R;
211     a += 9;
212     *R = a;
213   }
214 };
215 
216 struct SD {
217   int arr[32];
218   void foo() {
219     int a = *R;
220     a += 10;
221     *R = a;
222   }
223   SD() {
224     int a = *R;
225     a += 11;
226     *R = a;
227   }
228   ~SD() {
229     int a = *R;
230     #pragma omp target
231     a += 12;
232     *R = a;
233   }
234 };
235 
236 struct SE {
237   int arr[64];
238   void foo() {
239     int a = *R;
240     #pragma omp target if(0)
241     a += 13;
242     *R = a;
243   }
244   SE() {
245     int a = *R;
246     #pragma omp target
247     a += 14;
248     *R = a;
249   }
250   ~SE() {
251     int a = *R;
252     #pragma omp target
253     a += 15;
254     *R = a;
255   }
256 };
257 
258 template <int x>
259 struct ST {
260   int arr[128 + x];
261   void foo() {
262     int a = *R;
263     #pragma omp target
264     a += 16 + x;
265     *R = a;
266   }
267   ST() {
268     int a = *R;
269     #pragma omp target
270     a += 17 + x;
271     *R = a;
272   }
273   ~ST() {
274     int a = *R;
275     #pragma omp target
276     a += 18 + x;
277     *R = a;
278   }
279 };
280 
281 // We have to make sure we us all the target regions:
282 //CHECK-DAG: define internal void @[[NAME1]](
283 //CHECK-DAG: call void @[[NAME1]](
284 //CHECK-DAG: define internal void @[[NAME2]](
285 //CHECK-DAG: call void @[[NAME2]](
286 //CHECK-DAG: define internal void @[[NAME3]](
287 //CHECK-DAG: call void @[[NAME3]](
288 //CHECK-DAG: define internal void @[[NAME4]](
289 //CHECK-DAG: call void @[[NAME4]](
290 //CHECK-DAG: define internal void @[[NAME5]](
291 //CHECK-DAG: call void @[[NAME5]](
292 //CHECK-DAG: define internal void @[[NAME6]](
293 //CHECK-DAG: call void @[[NAME6]](
294 //CHECK-DAG: define internal void @[[NAME7]](
295 //CHECK-DAG: call void @[[NAME7]](
296 //CHECK-DAG: define internal void @[[NAME8]](
297 //CHECK-DAG: call void @[[NAME8]](
298 //CHECK-DAG: define internal void @[[NAME9]](
299 //CHECK-DAG: call void @[[NAME9]](
300 //CHECK-DAG: define internal void @[[NAME10]](
301 //CHECK-DAG: call void @[[NAME10]](
302 //CHECK-DAG: define internal void @[[NAME11]](
303 //CHECK-DAG: call void @[[NAME11]](
304 //CHECK-DAG: define internal void @[[NAME12]](
305 //CHECK-DAG: call void @[[NAME12]](
306 
307 //TCHECK-DAG: define weak{{.*}} void @[[NAME1]](
308 //TCHECK-DAG: define weak{{.*}} void @[[NAME2]](
309 //TCHECK-DAG: define weak{{.*}} void @[[NAME3]](
310 //TCHECK-DAG: define weak{{.*}} void @[[NAME4]](
311 //TCHECK-DAG: define weak{{.*}} void @[[NAME5]](
312 //TCHECK-DAG: define weak{{.*}} void @[[NAME6]](
313 //TCHECK-DAG: define weak{{.*}} void @[[NAME7]](
314 //TCHECK-DAG: define weak{{.*}} void @[[NAME8]](
315 //TCHECK-DAG: define weak{{.*}} void @[[NAME9]](
316 //TCHECK-DAG: define weak{{.*}} void @[[NAME10]](
317 //TCHECK-DAG: define weak{{.*}} void @[[NAME11]](
318 //TCHECK-DAG: define weak{{.*}} void @[[NAME12]](
319 
320 // CHECK-NTARGET-NOT: __tgt_target
321 
322 // TCHECK-NOT: __tgt_target
323 
324 // We have 2 initializers with priority 500
325 //CHECK: define internal void [[P500]](
326 //CHECK:     call void @{{.+}}()
327 //CHECK:     call void @{{.+}}()
328 //CHECK-NOT: call void @{{.+}}()
329 //CHECK:     ret void
330 
331 // We have 1 initializers with priority 501
332 //CHECK: define internal void [[P501]](
333 //CHECK:     call void @{{.+}}()
334 //CHECK-NOT: call void @{{.+}}()
335 //CHECK:     ret void
336 
337 // We have 6 initializers with default priority
338 //CHECK: define internal void [[PMAX]](
339 //CHECK:     call void @{{.+}}()
340 //CHECK:     call void @{{.+}}()
341 //CHECK:     call void @{{.+}}()
342 //CHECK:     call void @{{.+}}()
343 //CHECK:     call void @{{.+}}()
344 //CHECK:     call void @{{.+}}()
345 //CHECK-NOT: call void @{{.+}}()
346 //CHECK:     ret void
347 
348 static __attribute__((init_priority(500))) SA a1;
349 SA a2;
350 SB __attribute__((init_priority(500))) b1;
351 SB __attribute__((init_priority(501))) b2;
352 static SC c1;
353 SD d1;
354 SE e1;
355 ST<100> t1;
356 ST<1000> t2;
357 
358 
359 int bar(int a){
360   int r = a;
361 
362   a1.foo();
363   a2.foo();
364   b1.foo();
365   b2.foo();
366   c1.foo();
367   d1.foo();
368   e1.foo();
369   t1.foo();
370   t2.foo();
371 
372   #pragma omp target
373   ++r;
374 
375   return r + *R;
376 }
377 
378 // Check metadata is properly generated:
379 // CHECK:     !omp_offload.info = !{!{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}}
380 // CHECK-DAG: = !{i32 0, i32 [[DEVID:-?[0-9]+]], i32 [[FILEID:-?[0-9]+]], !"_ZN2SB3fooEv", i32 {{[0-9]+}}, i32 0, i32 {{[0-9]+}}}
381 // CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2SDD1Ev", i32 {{[0-9]+}}, i32 0, i32 {{[0-9]+}}}
382 // CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2SEC1Ev", i32 {{[0-9]+}}, i32 0, i32 {{[0-9]+}}}
383 // CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2SED1Ev", i32 {{[0-9]+}}, i32 0, i32 {{[0-9]+}}}
384 // CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi1000EE3fooEv", i32 {{[0-9]+}}, i32 0, i32 {{[0-9]+}}}
385 // CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi100EEC1Ev", i32 {{[0-9]+}}, i32 0, i32 {{[0-9]+}}}
386 // CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_Z3bari", i32 {{[0-9]+}}, i32 0, i32 {{[0-9]+}}}
387 // CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi100EED1Ev", i32 {{[0-9]+}}, i32 0, i32 {{[0-9]+}}}
388 // CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi1000EEC1Ev", i32 {{[0-9]+}}, i32 0, i32 {{[0-9]+}}}
389 // CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi1000EED1Ev", i32 {{[0-9]+}}, i32 0, i32 {{[0-9]+}}}
390 // CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi100EE3fooEv", i32 {{[0-9]+}}, i32 0, i32 {{[0-9]+}}}
391 // CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2SCC1Ev", i32 {{[0-9]+}}, i32 0, i32 {{[0-9]+}}}
392 
393 // TCHECK:     !omp_offload.info = !{!{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}}
394 // TCHECK-DAG: = !{i32 0, i32 [[DEVID:-?[0-9]+]], i32 [[FILEID:-?[0-9]+]], !"_ZN2SB3fooEv", i32 {{[0-9]+}}, i32 0, i32 {{[0-9]+}}}
395 // TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2SDD1Ev", i32 {{[0-9]+}}, i32 0, i32 {{[0-9]+}}}
396 // TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2SEC1Ev", i32 {{[0-9]+}}, i32 0, i32 {{[0-9]+}}}
397 // TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2SED1Ev", i32 {{[0-9]+}}, i32 0, i32 {{[0-9]+}}}
398 // TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi1000EE3fooEv", i32 {{[0-9]+}}, i32 0, i32 {{[0-9]+}}}
399 // TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi100EEC1Ev", i32 {{[0-9]+}}, i32 0, i32 {{[0-9]+}}}
400 // TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_Z3bari", i32 {{[0-9]+}}, i32 0, i32 {{[0-9]+}}}
401 // TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi100EED1Ev", i32 {{[0-9]+}}, i32 0, i32 {{[0-9]+}}}
402 // TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi1000EEC1Ev", i32 {{[0-9]+}}, i32 0, i32 {{[0-9]+}}}
403 // TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi1000EED1Ev", i32 {{[0-9]+}}, i32 0, i32 {{[0-9]+}}}
404 // TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi100EE3fooEv", i32 {{[0-9]+}}, i32 0, i32 {{[0-9]+}}}
405 // TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2SCC1Ev", i32 {{[0-9]+}}, i32 0, i32 {{[0-9]+}}}
406 
407 #endif
408