1 | // Test host codegen. |
2 | // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64 |
3 | // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -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 -fopenmp-version=45 -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 --check-prefix CHECK --check-prefix CHECK-64 |
5 | // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 |
6 | // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -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 -fopenmp-version=45 -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 --check-prefix CHECK --check-prefix CHECK-32 |
8 | |
9 | // RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -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 -fopenmp-version=45 -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 -fopenmp-version=45 -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 -fopenmp-version=45 -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 -fopenmp-version=45 -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 -fopenmp-version=45 -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 -fopenmp-version=45 -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 -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-64 |
20 | // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s |
21 | // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-64 |
22 | // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -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 -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-32 |
24 | // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s |
25 | // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-32 |
26 | |
27 | // RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -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 -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck --check-prefix SIMD-ONLY1 %s |
29 | // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s |
30 | // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-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 -fopenmp-version=45 -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 -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck --check-prefix SIMD-ONLY1 %s |
33 | // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s |
34 | // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-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 | // expected-no-diagnostics |
38 | #ifndef HEADER |
39 | #define HEADER |
40 | |
41 | // CHECK-DAG: %struct.ident_t = type { i32, i32, i32, i32, i8* } |
42 | // CHECK-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00" |
43 | // CHECK-DAG: [[DEF_LOC:@.+]] = private unnamed_addr global %struct.ident_t { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) } |
44 | |
45 | // CHECK-DAG: [[S1:%.+]] = type { double } |
46 | // CHECK-DAG: [[ENTTY:%.+]] = type { i8*, i8*, i[[SZ:32|64]], i32, i32 } |
47 | // CHECK-DAG: [[DEVTY:%.+]] = type { i8*, i8*, [[ENTTY]]*, [[ENTTY]]* } |
48 | // CHECK-DAG: [[DSCTY:%.+]] = type { i32, [[DEVTY]]*, [[ENTTY]]*, [[ENTTY]]* } |
49 | |
50 | // TCHECK: [[ENTTY:%.+]] = type { i8*, i8*, i{{32|64}}, i32, i32 } |
51 | |
52 | // CHECK-DAG: $[[REGFN:\.omp_offloading\..+]] = comdat |
53 | |
54 | // We have 6 target regions |
55 | |
56 | // CHECK-DAG: @{{.*}} = weak constant i8 0 |
57 | // CHECK-DAG: @{{.*}} = weak constant i8 0 |
58 | // CHECK-DAG: @{{.*}} = weak constant i8 0 |
59 | // CHECK-DAG: @{{.*}} = weak constant i8 0 |
60 | // CHECK-DAG: @{{.*}} = weak constant i8 0 |
61 | // CHECK-DAG: @{{.*}} = weak constant i8 0 |
62 | |
63 | // TCHECK: @{{.+}} = weak constant [[ENTTY]] |
64 | // TCHECK: @{{.+}} = weak constant [[ENTTY]] |
65 | // TCHECK: @{{.+}} = weak constant [[ENTTY]] |
66 | // TCHECK: @{{.+}} = weak constant [[ENTTY]] |
67 | // TCHECK: @{{.+}} = weak constant [[ENTTY]] |
68 | // TCHECK: @{{.+}} = weak constant [[ENTTY]] |
69 | |
70 | // Check if offloading descriptor is created. |
71 | // CHECK: [[ENTBEGIN:@.+]] = external constant [[ENTTY]] |
72 | // CHECK: [[ENTEND:@.+]] = external constant [[ENTTY]] |
73 | // CHECK: [[DEVBEGIN:@.+]] = extern_weak constant i8 |
74 | // CHECK: [[DEVEND:@.+]] = extern_weak constant i8 |
75 | // CHECK: [[IMAGES:@.+]] = internal unnamed_addr constant [1 x [[DEVTY]]] [{{.+}} { i8* [[DEVBEGIN]], i8* [[DEVEND]], [[ENTTY]]* [[ENTBEGIN]], [[ENTTY]]* [[ENTEND]] }], comdat($[[REGFN]]) |
76 | // CHECK: [[DESC:@.+]] = internal constant [[DSCTY]] { i32 1, [[DEVTY]]* getelementptr inbounds ([1 x [[DEVTY]]], [1 x [[DEVTY]]]* [[IMAGES]], i32 0, i32 0), [[ENTTY]]* [[ENTBEGIN]], [[ENTTY]]* [[ENTEND]] }, comdat($[[REGFN]]) |
77 | |
78 | // Check target registration is registered as a Ctor. |
79 | // CHECK: appending global [1 x { i32, void ()*, i8* }] [{ i32, void ()*, i8* } { i32 0, void ()* @[[REGFN]], i8* bitcast (void ()* @[[REGFN]] to i8*) }] |
80 | |
81 | |
82 | template<typename tx> |
83 | tx ftemplate(int n) { |
84 | tx a = 0; |
85 | |
86 | #pragma omp target teams num_teams(tx(20)) |
87 | { |
88 | } |
89 | |
90 | short b = 1; |
91 | #pragma omp target teams num_teams(b) |
92 | { |
93 | a += b; |
94 | } |
95 | |
96 | return a; |
97 | } |
98 | |
99 | static |
100 | int fstatic(int n) { |
101 | |
102 | #pragma omp target teams num_teams(n) |
103 | { |
104 | } |
105 | |
106 | #pragma omp target teams num_teams(32+n) |
107 | { |
108 | } |
109 | |
110 | return n+1; |
111 | } |
112 | |
113 | struct S1 { |
114 | double a; |
115 | |
116 | int r1(int n){ |
117 | int b = 1; |
118 | |
119 | #pragma omp target teams num_teams(n-b) |
120 | { |
121 | this->a = (double)b + 1.5; |
122 | } |
123 | |
124 | #pragma omp target teams num_teams(1024) |
125 | { |
126 | this->a = 2.5; |
127 | } |
128 | |
129 | return (int)a; |
130 | } |
131 | }; |
132 | |
133 | // CHECK: define {{.*}}@{{.*}}bar{{.*}} |
134 | int bar(int n){ |
135 | int a = 0; |
136 | |
137 | S1 S; |
138 | // CHECK: call {{.*}}i32 [[FS1:@.+]]([[S1]]* {{.*}}, i32 {{.*}}) |
139 | a += S.r1(n); |
140 | |
141 | // CHECK: call {{.*}}i32 [[FSTATIC:@.+]](i32 {{.*}}) |
142 | a += fstatic(n); |
143 | |
144 | // CHECK: call {{.*}}i32 [[FTEMPLATE:@.+]](i32 {{.*}}) |
145 | a += ftemplate<int>(n); |
146 | |
147 | return a; |
148 | } |
149 | |
150 | |
151 | |
152 | // |
153 | // CHECK: define {{.*}}[[FS1]]([[S1]]* {{%.+}}, i32 {{[^%]*}}[[PARM:%.+]]) |
154 | // |
155 | // CHECK-DAG: store i32 [[PARM]], i32* [[N_ADDR:%.+]], align |
156 | // CHECK: store i32 1, i32* [[B:%.+]], align |
157 | // CHECK: [[NV:%.+]] = load i32, i32* [[N_ADDR]], align |
158 | // CHECK: [[BV:%.+]] = load i32, i32* [[B]], align |
159 | // CHECK: [[SUB:%.+]] = sub nsw i32 [[NV]], [[BV]] |
160 | // CHECK: store i32 [[SUB]], i32* [[CAPE_ADDR:%.+]], align |
161 | // CHECK: [[CEV:%.+]] = load i32, i32* [[CAPE_ADDR]], align |
162 | // CHECK-64: [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPEC_ADDR:%.+]] to i32* |
163 | // CHECK-64: store i32 [[CEV]], i32* [[CONV]], align |
164 | // CHECK-32: store i32 [[CEV]], i32* [[CAPEC_ADDR:%.+]], align |
165 | // CHECK: [[ARG:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR]], align |
166 | // CHECK: [[TEAMS:%.+]] = load i32, i32* [[CAPE_ADDR]], align |
167 | // |
168 | // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 4, {{.*}}, i32 [[TEAMS]], i32 0) |
169 | // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 |
170 | // CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] |
171 | // |
172 | // CHECK: [[FAIL]] |
173 | // CHECK: call void [[HVT1:@.+]]([[S1]]* {{%.+}}, i[[SZ]] {{%.+}}, i[[SZ]] [[ARG]]) |
174 | // CHECK: br label {{%?}}[[END]] |
175 | // CHECK: [[END]] |
176 | // |
177 | // |
178 | // |
179 | // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 2, {{.+}}, i32 1024, i32 0) |
180 | // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 |
181 | // CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] |
182 | // |
183 | // CHECK: [[FAIL]] |
184 | // CHECK: call void [[HVT2:@.+]]([[S1]]* {{[^,]+}}) |
185 | // CHECK: br label {{%?}}[[END]] |
186 | // CHECK: [[END]] |
187 | // |
188 | |
189 | |
190 | |
191 | |
192 | |
193 | |
194 | // |
195 | // CHECK: define {{.*}}[[FSTATIC]](i32 {{[^%]*}}[[PARM:%.+]]) |
196 | // |
197 | // CHECK-DAG: store i32 [[PARM]], i32* [[N_ADDR:%.+]], align |
198 | // CHECK: [[NV:%.+]] = load i32, i32* [[N_ADDR]], align |
199 | // CHECK: store i32 [[NV]], i32* [[CAPE_ADDR:%.+]], align |
200 | // CHECK: [[CEV:%.+]] = load i32, i32* [[CAPE_ADDR]], align |
201 | // CHECK-64: [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPEC_ADDR:%.+]] to i32* |
202 | // CHECK-64: store i32 [[CEV]], i32* [[CONV]], align |
203 | // CHECK-32: store i32 [[CEV]], i32* [[CAPEC_ADDR:%.+]], align |
204 | // CHECK: [[ARG:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR]], align |
205 | // CHECK: [[TEAMS:%.+]] = load i32, i32* [[CAPE_ADDR]], align |
206 | // |
207 | // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 1, {{.*}}, i32 [[TEAMS]], i32 0) |
208 | // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 |
209 | // CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] |
210 | // |
211 | // CHECK: [[FAIL]] |
212 | // CHECK: call void [[HVT3:@.+]](i[[SZ]] [[ARG]]) |
213 | // CHECK: br label {{%?}}[[END]] |
214 | // CHECK: [[END]] |
215 | // |
216 | // |
217 | // |
218 | // CHECK: [[NV:%.+]] = load i32, i32* [[N_ADDR]], align |
219 | // CHECK: [[ADD:%.+]] = add nsw i32 32, [[NV]] |
220 | // CHECK: store i32 [[ADD]], i32* [[CAPE_ADDR:%.+]], align |
221 | // CHECK: [[CEV:%.+]] = load i32, i32* [[CAPE_ADDR]], align |
222 | // CHECK-64: [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPEC_ADDR:%.+]] to i32* |
223 | // CHECK-64: store i32 [[CEV]], i32* [[CONV]], align |
224 | // CHECK-32: store i32 [[CEV]], i32* [[CAPEC_ADDR:%.+]], align |
225 | // CHECK: [[ARG:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR]], align |
226 | // CHECK: [[TEAMS:%.+]] = load i32, i32* [[CAPE_ADDR]], align |
227 | // |
228 | // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 1, {{.*}}, i32 [[TEAMS]], i32 0) |
229 | // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 |
230 | // CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] |
231 | // |
232 | // CHECK: [[FAIL]] |
233 | // CHECK: call void [[HVT4:@.+]](i[[SZ]] [[ARG]]) |
234 | // CHECK: br label {{%?}}[[END]] |
235 | // CHECK: [[END]] |
236 | // |
237 | |
238 | |
239 | |
240 | |
241 | |
242 | |
243 | // |
244 | // CHECK: define {{.*}}[[FTEMPLATE]] |
245 | // |
246 | // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 0, {{.*}}, i32 20, i32 0) |
247 | // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 |
248 | // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] |
249 | // |
250 | // CHECK: [[FAIL]] |
251 | // CHECK: call void [[HVT5:@.+]]() |
252 | // CHECK: br label {{%?}}[[END]] |
253 | // |
254 | // CHECK: [[END]] |
255 | // |
256 | // |
257 | // |
258 | // CHECK: store i16 1, i16* [[B:%.+]], align |
259 | // CHECK: [[BV:%.+]] = load i16, i16* [[B]], align |
260 | // CHECK: store i16 [[BV]], i16* [[CAPE_ADDR:%.+]], align |
261 | // CHECK: [[CEV:%.+]] = load i16, i16* [[CAPE_ADDR]], align |
262 | // CHECK: [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPEC_ADDR:%.+]] to i16* |
263 | // CHECK: store i16 [[CEV]], i16* [[CONV]], align |
264 | // CHECK: [[ARG:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR]], align |
265 | // CHECK: [[T:%.+]] = load i16, i16* [[CAPE_ADDR]], align |
266 | // CHECK: [[TEAMS:%.+]] = sext i16 [[T]] to i32 |
267 | // |
268 | // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 3, {{.*}}, i32 [[TEAMS]], i32 0) |
269 | // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 |
270 | // CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] |
271 | // |
272 | // CHECK: [[FAIL]] |
273 | // CHECK: call void [[HVT6:@.+]](i[[SZ]] {{%.+}}, i[[SZ]] {{%.+}}, i[[SZ]] [[ARG]]) |
274 | // CHECK: br label {{%?}}[[END]] |
275 | // CHECK: [[END]] |
276 | // |
277 | |
278 | |
279 | |
280 | |
281 | |
282 | |
283 | // Check that the offloading functions are emitted and that the parallel function |
284 | // is appropriately guarded. |
285 | |
286 | // CHECK: define internal void [[HVT1]]([[S1]]* {{%.+}}, i[[SZ]] [[PARM1:%.+]], i[[SZ]] [[PARM2:%.+]]) |
287 | // CHECK-DAG: store i[[SZ]] [[PARM2]], i[[SZ]]* [[CAPE_ADDR:%.+]], align |
288 | // CHECK-64: [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPE_ADDR]] to i32* |
289 | // CHECK-64: [[NT:%.+]] = load i32, i32* [[CONV]], align |
290 | // CHECK-32: [[NT:%.+]] = load i32, i32* [[CAPE_ADDR]], align |
291 | // CHECK: call i32 @__kmpc_push_num_teams(%struct.ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 [[NT]], i32 0) |
292 | // CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* [[DEF_LOC]], i32 2, |
293 | // |
294 | // |
295 | |
296 | |
297 | // CHECK: define internal void [[HVT2]]([[S1]]* {{%.+}}) |
298 | // CHECK: call i32 @__kmpc_push_num_teams(%struct.ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 1024, i32 0) |
299 | // CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* [[DEF_LOC]], i32 1, |
300 | // |
301 | // |
302 | |
303 | |
304 | |
305 | |
306 | |
307 | |
308 | |
309 | |
310 | // CHECK: define internal void [[HVT3]](i[[SZ]] [[PARM:%.+]]) |
311 | // CHECK-DAG: store i[[SZ]] [[PARM]], i[[SZ]]* [[CAPE_ADDR:%.+]], align |
312 | // CHECK-64: [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPE_ADDR]] to i32* |
313 | // CHECK-64: [[NT:%.+]] = load i32, i32* [[CONV]], align |
314 | // CHECK-32: [[NT:%.+]] = load i32, i32* [[CAPE_ADDR]], align |
315 | // CHECK: call i32 @__kmpc_push_num_teams(%struct.ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 [[NT]], i32 0) |
316 | // CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* [[DEF_LOC]], i32 0, |
317 | // |
318 | // |
319 | // CHECK: define internal void [[HVT4]](i[[SZ]] [[PARM:%.+]]) |
320 | // CHECK-DAG: store i[[SZ]] [[PARM]], i[[SZ]]* [[CAPE_ADDR:%.+]], align |
321 | // CHECK-64: [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPE_ADDR]] to i32* |
322 | // CHECK-64: [[NT:%.+]] = load i32, i32* [[CONV]], align |
323 | // CHECK-32: [[NT:%.+]] = load i32, i32* [[CAPE_ADDR]], align |
324 | // CHECK: call i32 @__kmpc_push_num_teams(%struct.ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 [[NT]], i32 0) |
325 | // CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* [[DEF_LOC]], i32 0, |
326 | // |
327 | // |
328 | |
329 | |
330 | |
331 | |
332 | |
333 | // CHECK: define internal void [[HVT5]]( |
334 | // CHECK: call i32 @__kmpc_push_num_teams(%struct.ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 20, i32 0) |
335 | // CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* [[DEF_LOC]], i32 0, |
336 | // |
337 | // |
338 | |
339 | |
340 | // CHECK: define internal void [[HVT6]](i[[SZ]] [[PARM1:%.+]], i[[SZ]] [[PARM2:%.+]], i[[SZ]] [[PARM3:%.+]]) |
341 | // CHECK-DAG: store i[[SZ]] [[PARM3]], i[[SZ]]* [[CAPE_ADDR:%.+]], align |
342 | // CHECK: [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPE_ADDR]] to i16* |
343 | // CHECK: [[T:%.+]] = load i16, i16* [[CONV]], align |
344 | // CHECK: [[NT:%.+]] = sext i16 [[T]] to i32 |
345 | // CHECK: call i32 @__kmpc_push_num_teams(%struct.ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 [[NT]], i32 0) |
346 | // CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* [[DEF_LOC]], i32 2, |
347 | // |
348 | // |
349 | |
350 | |
351 | |
352 | #endif |
353 | |