1 // expected-no-diagnostics
2 #ifndef HEADER
3 #define HEADER
4 // Test host codegen.
5 // RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fomptargets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-64
6 // RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fomptargets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
7 // RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fomptargets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-64
8 // RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fomptargets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-32
9 // RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fomptargets=i386-pc-linux-gnu -emit-pch -o %t %s
10 // RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -triple i386-unknown-unknown -fomptargets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-32
11 #ifdef CK1
12 
13 int Gbla;
14 long long Gblb;
15 int &Gblc = Gbla;
16 
17 // CK1-LABEL: teams_argument_global_local
18 int teams_argument_global_local(int a){
19   int comp = 1;
20 
21   int la = 23;
22   float lc = 25.0;
23 
24   // CK1: call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 1, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0)
25   // CK1: call void @{{.+}}(i{{64|32}} %{{.+}})
26   #pragma omp target
27   #pragma omp teams
28   {
29     ++comp;
30   }
31 
32   // CK1-DAG: call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 2, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 [[NT:%[^,]+]], i32 0)
33   // CK1-DAG: [[NT]] = load i32, i32* [[NTA:%[^,]+]],
34 
35   // CK1: call void @{{.+}}(i{{64|32}} %{{.+}})
36   #pragma omp target
37   #pragma omp teams num_teams(la)
38   {
39     ++comp;
40   }
41 
42   // CK1-DAG: call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 2, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 [[NT:%[^,]+]])
43   // CK1-DAG: [[NT]] = load i32, i32* [[NTA:%[^,]+]],
44 
45   // CK1: call void @{{.+}}(i{{64|32}} %{{.+}})
46   #pragma omp target
47   #pragma omp teams thread_limit(la)
48   {
49     ++comp;
50   }
51 
52   // CK1-DAG: call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 5, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 [[NT:%[^,]+]], i32 [[TL:%[^,]+]])
53 
54   // CK1-DAG: [[NT]] = add nsw i32 [[NTA:%[^,]+]], [[NTB:%[^,]+]]
55   // CK1-DAG: [[NTA]] = load i32, i32* @Gbla,
56   // CK1-DAG: [[NTB]] = load i32, i32* %{{.+}},
57 
58   // CK1-DAG: [[TL]] = trunc i64 [[TLA:%[^,]+]] to i32
59   // CK1-DAG: [[TLA]] = add nsw i64 [[TLB:%[^,]+]], [[TLC:%[^,]+]]
60   // CK1-DAG: [[TLC]] = fptosi float [[TLD:%[^,]+]] to i64
61   // CK1-DAG: [[TLD]] = load float, float* %{{.+}},
62   // CK1-DAG: [[TLB]] = load i64, i64* @Gblb,
63 
64   // CK1: call void @{{.+}}(i{{.+}} {{.+}}, i{{.+}} {{.+}}, i{{.+}} {{.+}}, i{{.+}} {{.+}}, i{{.+}} {{.+}})
65   #pragma omp target
66   #pragma omp teams num_teams(Gbla+a) thread_limit(Gblb+(long long)lc)
67   {
68     ++comp;
69   }
70 
71   // CK1-DAG: call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 {{.+}}, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 [[NT:%[^,]+]], i32 [[TL:%[^,]+]])
72 
73   // CK1-DAG: [[NT]] = add nsw i32 [[NTA:%[^,]+]], 1
74   // CK1-DAG: [[NTA]] = load i32, i32* @Gbla,
75 
76   // CK1-DAG: [[TL]] = add nsw i32 [[TLA:%[^,]+]], 2
77   // CK1-DAG: [[TLA]] = load i32, i32* @Gbla,
78 
79   // CK1: call void @{{.+}}(i{{.+}} {{.+}}
80   #pragma omp target
81   #pragma omp teams num_teams(Gblc+1) thread_limit(Gblc+2)
82   {
83     comp += Gblc;
84   }
85 
86   return comp;
87 }
88 
89 #endif // CK1
90 
91 // Test host codegen.
92 // RUN: %clang_cc1 -DCK2 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fomptargets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-64
93 // RUN: %clang_cc1 -DCK2 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fomptargets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
94 // RUN: %clang_cc1 -DCK2 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fomptargets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-64
95 // RUN: %clang_cc1 -DCK2 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fomptargets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-32
96 // RUN: %clang_cc1 -DCK2 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fomptargets=i386-pc-linux-gnu -emit-pch -o %t %s
97 // RUN: %clang_cc1 -DCK2 -fopenmp -x c++ -triple i386-unknown-unknown -fomptargets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-32
98 #ifdef CK2
99 
100 // CK2-DAG: [[SSI:%.+]] = type { i32, float }
101 // CK2-DAG: [[SSL:%.+]] = type { i64, float }
102 template <typename T>
103 struct SS{
104   T a;
105   float b;
106 };
107 
108 SS<int> Gbla;
109 SS<long long> Gblb;
110 
111 // CK2-LABEL: teams_template_arg
112 int teams_template_arg(void) {
113   int comp = 1;
114 
115   SS<int> la;
116   SS<long long> lb;
117 
118   // CK2-DAG: call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 3, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 [[NT:%[^,]+]], i32 [[TL:%[^,]+]])
119 
120   // CK2-DAG: [[NT]] = load i32, i32* getelementptr inbounds ([[SSI]], [[SSI]]* @Gbla, i32 0, i32 0)
121 
122   // CK2-DAG: [[TL]] = trunc i64 [[TLA:%[^,]+]] to i32
123   // CK2-DAG: [[TLA]] = fptosi float [[TLB:%[^,]+]] to i64
124   // CK2-DAG: [[TLB]] = load float, float* [[TLC:%[^,]+]],
125   // CK2-DAG: [[TLC]] = getelementptr inbounds [[SSI]], [[SSI]]* %{{.+}}, i32 0, i32 1
126 
127   // CK2: call void @{{.+}}({{.+}} {{.+}}, {{.+}} {{.+}}, {{.+}} {{.+}})
128   #pragma omp target
129   #pragma omp teams num_teams(Gbla.a) thread_limit((long long)la.b)
130   {
131     ++comp;
132   }
133 
134   // CK2-DAG: call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 3, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 [[NT:%[^,]+]], i32 [[TL:%[^,]+]])
135 
136   // CK2-DAG: [[TL]] = trunc i64 [[TLD:%[^,]+]] to i32
137   // CK2-DAG: [[TLD]] = load i64, i64* getelementptr inbounds ([[SSL]], [[SSL]]* @Gblb, i32 0, i32 0),
138 
139   // CK2-DAG: [[NT]] = trunc i64 [[NTA:%[^,]+]] to i32
140   // CK2-DAG: [[NTA]] = fptosi float [[NTB:%[^,]+]] to i64
141   // CK2-DAG: [[NTB]] = load float, float* [[NTC:%[^,]+]],
142   // CK2-DAG: [[NTC]] = getelementptr inbounds [[SSL]], [[SSL]]* %{{.+}}, i32 0, i32 1
143 
144   // CK2: call void @{{.+}}({{.+}} {{.+}}, {{.+}} {{.+}}, {{.+}} {{.+}})
145   #pragma omp target
146   #pragma omp teams num_teams((long long)lb.b) thread_limit(Gblb.a)
147   {
148     ++comp;
149   }
150   return comp;
151 }
152 #endif // CK2
153 
154 // Test host codegen.
155 // RUN: %clang_cc1 -DCK3 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fomptargets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CK3 --check-prefix CK3-64
156 // RUN: %clang_cc1 -DCK3 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fomptargets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
157 // RUN: %clang_cc1 -DCK3 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fomptargets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK3 --check-prefix CK3-64
158 // RUN: %clang_cc1 -DCK3 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fomptargets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CK3 --check-prefix CK3-32
159 // RUN: %clang_cc1 -DCK3 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fomptargets=i386-pc-linux-gnu -emit-pch -o %t %s
160 // RUN: %clang_cc1 -DCK3 -fopenmp -x c++ -triple i386-unknown-unknown -fomptargets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK3 --check-prefix CK3-32
161 #ifdef CK3
162 
163 // CK3: [[SSI:%.+]] = type { i32, float }
164 // CK3-LABEL: teams_template_struct
165 
166 template <typename T, int X, long long Y>
167 struct SS{
168   T a;
169   float b;
170 
171   int foo(void) {
172     int comp = 1;
173 
174     // CK3-DAG: call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 2, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 [[NT:%[^,]+]], i32 123)
175 
176     // CK3-DAG: [[NT]] = load i32, i32* [[NTA:%[^,]+]],
177     // CK3-DAG: [[NTA]] = getelementptr inbounds [[SSI]], [[SSI]]* [[NTB:%[^,]+]], i32 0, i32 0
178     // CK3-DAG: [[NTB]] = load [[SSI]]*, [[SSI]]** %{{.+}},
179 
180     // CK3: call void @{{.+}}({{.+}} {{.+}}, {{.+}} {{.+}})
181     #pragma omp target
182     #pragma omp teams num_teams(a) thread_limit(X)
183     {
184       ++comp;
185     }
186 
187     // CK3-DAG: call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 2, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 456, i32 [[TL:%[^,]+]])
188 
189     // CK3-DAG: [[TL]] = add nsw i32 [[TLA:%[^,]+]], 123
190     // CK3-DAG: [[TLA]] = fptosi float [[TLB:%[^,]+]] to i32
191     // CK3-DAG: [[TLB]] = load float, float* [[TLC:%[^,]+]],
192     // CK3-DAG: [[TLC]] = getelementptr inbounds [[SSI]], [[SSI]]* [[THIS:%[^,]+]], i32 0, i32 1
193 
194     // CK3: call void @{{.+}}({{.+}} {{.+}}, {{.+}} {{.+}})
195     #pragma omp target
196     #pragma omp teams num_teams(Y) thread_limit((int)b+X)
197     {
198       ++comp;
199     }
200     return comp;
201   }
202 };
203 
204 int teams_template_struct(void) {
205   SS<int, 123, 456> V;
206   return V.foo();
207 
208 }
209 #endif // CK3
210 
211 // Test target codegen - host bc file has to be created first.
212 // RUN: %clang_cc1 -DCK4 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fomptargets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc
213 // RUN: %clang_cc1 -DCK4 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fomptargets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fomp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix CK4 --check-prefix CK4-64
214 // RUN: %clang_cc1 -DCK4 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fomptargets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-device -fomp-host-ir-file-path %t-ppc-host.bc -o %t %s
215 // RUN: %clang_cc1 -DCK4 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fomptargets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-device -fomp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK4 --check-prefix CK4-64
216 // RUN: %clang_cc1 -DCK4 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fomptargets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc
217 // RUN: %clang_cc1 -DCK4 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fomptargets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-device -fomp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CK4 --check-prefix CK4-32
218 // RUN: %clang_cc1 -DCK4 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fomptargets=i386-pc-linux-gnu -emit-pch -fopenmp-is-device -fomp-host-ir-file-path %t-x86-host.bc -o %t %s
219 // RUN: %clang_cc1 -DCK4 -fopenmp -x c++ -triple i386-unknown-unknown -fomptargets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-device -fomp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK4 --check-prefix CK4-32
220 
221 #ifdef CK4
222 
223 // CK4-DAG: %ident_t = type { i32, i32, i32, i32, i8* }
224 // CK4-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
225 // CK4-DAG: [[DEF_LOC_0:@.+]] = private unnamed_addr constant %ident_t { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) }
226 // CK4-DEBUG-DAG: [[LOC1:@.+]] = private unnamed_addr constant [{{.+}} x i8] c";{{.*}}teams_codegen.cpp;main;[[@LINE+14]];9;;\00"
227 // CK4-DEBUG-DAG: [[LOC2:@.+]] = private unnamed_addr constant [{{.+}} x i8] c";{{.*}}teams_codegen.cpp;tmain;[[@LINE+7]];9;;\00"
228 
229 template <typename T>
230 int tmain(T argc) {
231 #pragma omp target
232 #pragma omp teams
233   argc = 0;
234   return 0;
235 }
236 
237 int main (int argc, char **argv) {
238 #pragma omp target
239 #pragma omp teams
240   argc = 0;
241   return tmain(argv);
242 }
243 
244 // CK4:  define {{.*}}void @{{[^,]+}}(i{{.+}} %[[ARGC:.+]])
245 // CK4:  [[ARGCADDR:%.+]] = alloca i{{.+}}
246 // CK4:  store i{{.+}} %[[ARGC]], i{{.+}}* [[ARGCADDR]]
247 // CK4-64:  [[CONV:%.+]] = bitcast i64* [[ARGCADDR]] to i32*
248 // CK4-64:  call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC_0]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* {{.+}} to void (i32*, i32*, ...)*), i32* [[CONV]])
249 // CK4-32:  call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC_0]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* {{.+}} to void (i32*, i32*, ...)*), i32* [[ARGCADDR]])
250 // CK4:  ret void
251 // CK4-NEXT: }
252 
253 // CK4:  define {{.*}}void @{{[^,]+}}(i8*** dereferenceable({{.}}) [[ARGC1:%.+]])
254 // CK4:  [[ARGCADDR1:%.+]] = alloca i8***
255 // CK4:  store i8*** [[ARGC1]], i8**** [[ARGCADDR1]]
256 // CK4:  [[CONV1:%.+]] = load i8***, i8**** [[ARGCADDR1]]
257 // CK4:  call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC_0]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i8***)* {{.+}} to void (i32*, i32*, ...)*), i8*** [[CONV1]])
258 
259 
260 #endif // CK4
261 
262 // Test target codegen - host bc file has to be created first.
263 // RUN: %clang_cc1 -DCK5 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fomptargets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc
264 // RUN: %clang_cc1 -DCK5 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fomptargets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fomp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix CK5 --check-prefix CK5-64
265 // RUN: %clang_cc1 -DCK5 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fomptargets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-device -fomp-host-ir-file-path %t-ppc-host.bc -o %t %s
266 // RUN: %clang_cc1 -DCK5 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fomptargets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-device -fomp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK5 --check-prefix CK5-64
267 // RUN: %clang_cc1 -DCK5 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fomptargets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc
268 // RUN: %clang_cc1 -DCK5 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fomptargets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-device -fomp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CK5 --check-prefix CK5-32
269 // RUN: %clang_cc1 -DCK5 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fomptargets=i386-pc-linux-gnu -emit-pch -fopenmp-is-device -fomp-host-ir-file-path %t-x86-host.bc -o %t %s
270 // RUN: %clang_cc1 -DCK5 -fopenmp -x c++ -triple i386-unknown-unknown -fomptargets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-device -fomp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK5 --check-prefix CK5-32
271 
272 // expected-no-diagnostics
273 #ifdef CK5
274 
275 // CK5-DAG: %ident_t = type { i32, i32, i32, i32, i8* }
276 // CK5-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
277 // CK5-DAG: [[DEF_LOC_0:@.+]] = private unnamed_addr constant %ident_t { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) }
278 // CK5-DEBUG-DAG: [[LOC1:@.+]] = private unnamed_addr constant [{{.+}} x i8] c";{{.*}}teams_codegen.cpp;main;[[@LINE+14]];9;;\00"
279 // CK5-DEBUG-DAG: [[LOC2:@.+]] = private unnamed_addr constant [{{.+}} x i8] c";{{.*}}teams_codegen.cpp;tmain;[[@LINE+7]];9;;\00"
280 
281 template <typename T>
282 int tmain(T argc) {
283   int a = 10;
284   int b = 5;
285 #pragma omp target
286 #pragma omp teams num_teams(a) thread_limit(b)
287   {
288   argc = 0;
289   }
290   return 0;
291 }
292 
293 int main (int argc, char **argv) {
294   int a = 20;
295   int b = 5;
296 #pragma omp target
297 #pragma omp teams num_teams(a) thread_limit(b)
298   {
299   argc = 0;
300   }
301   return tmain(argv);
302 }
303 
304 // CK5:  define {{.*}}void @{{[^,]+}}(i{{.+}} [[AP:%.+]], i{{.+}} [[BP:%.+]], i{{.+}} [[ARGC:.+]])
305 // CK5:  [[AADDR:%.+]] = alloca i{{.+}}
306 // CK5:  [[BADDR:%.+]] = alloca i{{.+}}
307 // CK5:  [[ARGCADDR:%.+]] = alloca i{{.+}}
308 // CK5:  [[GBL_TH_NUM:%.+]] = call i32 @__kmpc_global_thread_num(%ident_t* [[DEF_LOC_0]])
309 // CK5:  store i{{.+}} [[AP]], i{{.+}}* [[AADDR]]
310 // CK5:  store i{{.+}} [[BP]], i{{.+}}* [[BADDR]]
311 // CK5:  store i{{.+}} [[ARGC]], i{{.+}}* [[ARGCADDR]]
312 // CK5-64:  [[ACONV:%.+]] = bitcast i64* [[AADDR]] to i32*
313 // CK5-64:  [[BCONV:%.+]] = bitcast i64* [[BADDR]] to i32*
314 // CK5-64:  [[CONV:%.+]] = bitcast i64* [[ARGCADDR]] to i32*
315 // CK5-64:  [[ACONVVAL:%.+]] = load i32, i32* [[ACONV]]
316 // CK5-64:  [[BCONVVAL:%.+]] = load i32, i32* [[BCONV]]
317 // CK5-32:  [[ACONVVAL:%.+]] = load i32, i32* [[AADDR]]
318 // CK5-32:  [[BCONVVAL:%.+]] = load i32, i32* [[BADDR]]
319 // CK5:  {{.+}} = call i32 @__kmpc_push_num_teams(%ident_t* [[DEF_LOC_0]], i32 [[GBL_TH_NUM]], i32 [[ACONVVAL]], i32 [[BCONVVAL]])
320 // CK5-64:  call void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC_0]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined. to void (i32*, i32*, ...)*), i32* [[CONV]])
321 // CK5-32:  call void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC_0]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined. to void (i32*, i32*, ...)*), i32* [[ARGCADDR]])
322 
323 // CK5:  define {{.*}}void @{{[^,]+}}(i{{.+}} dereferenceable({{.+}}) [[AP:%.+]], i{{.+}} dereferenceable({{.+}}) [[BP:%.+]], i{{.+}} dereferenceable({{.+}}) [[ARGC:%.+]])
324 // CK5:  [[AADDR:%.+]] = alloca i{{.+}}
325 // CK5:  [[BADDR:%.+]] = alloca i{{.+}}
326 // CK5:  [[ARGCADDR:%.+]] = alloca i{{.+}}
327 // CK5:  [[GBL_TH_NUM:%.+]] = call i32 @__kmpc_global_thread_num(%ident_t* [[DEF_LOC_0]])
328 // CK5:  store i{{.+}} [[AP]], i{{.+}}* [[AADDR]]
329 // CK5:  store i{{.+}} [[BP]], i{{.+}}* [[BADDR]]
330 // CK5:  store i{{.+}} [[ARGC]], i{{.+}}* [[ARGCADDR]]
331 // CK5:  [[A_ADDR_VAL:%.+]] = load i32*, i32** [[AADDR]]
332 // CK5:  [[B_ADDR_VAL:%.+]] = load i32*, i32** [[BADDR]]
333 // CK5:  [[ARGC_ADDR_VAL:%.+]] = load i{{.+}}, i{{.+}}* [[ARGCADDR]]
334 // CK5:  [[A_VAL:%.+]] = load i32, i32* [[A_ADDR_VAL]]
335 // CK5:  [[B_VAL:%.+]] = load i32, i32* [[B_ADDR_VAL]]
336 // CK5:  {{.+}} = call i32 @__kmpc_push_num_teams(%ident_t* [[DEF_LOC_0]], i32 [[GBL_TH_NUM]], i32 [[A_VAL]], i32 [[B_VAL]])
337 // CK5:  call void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC_0]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i{{.+}})* @.omp_outlined.{{.+}} to void (i32*, i32*, ...)*), i{{.+}} [[ARGC_ADDR_VAL]])
338 // CK5:  ret void
339 // CK5-NEXT: }
340 
341 #endif // CK5
342 #endif
343