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 thread_limit(tx(20)) |
87 | { |
88 | } |
89 | |
90 | short b = 1; |
91 | #pragma omp target teams num_teams(b) thread_limit(1024) |
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) thread_limit(n*32) |
103 | { |
104 | } |
105 | |
106 | #pragma omp target teams thread_limit(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 thread_limit(n-b) |
120 | { |
121 | this->a = (double)b + 1.5; |
122 | } |
123 | |
124 | #pragma omp target teams thread_limit(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: [[TL:%.+]] = load i32, i32* [[CAPE_ADDR]], align |
167 | // |
168 | // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 4, {{.*}}, i32 0, i32 [[TL]]) |
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 0, i32 1024) |
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_ADDR1:%.+]], align |
200 | // CHECK: [[NV:%.+]] = load i32, i32* [[N_ADDR]], align |
201 | // CHECK: [[MUL:%.+]] = mul nsw i32 [[NV]], 32 |
202 | // CHECK: store i32 [[MUL]], i32* [[CAPE_ADDR2:%.+]], align |
203 | // CHECK: [[CEV:%.+]] = load i32, i32* [[CAPE_ADDR1]], align |
204 | // CHECK-64: [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPEC_ADDR1:%.+]] to i32* |
205 | // CHECK-64: store i32 [[CEV]], i32* [[CONV]], align |
206 | // CHECK-32: store i32 [[CEV]], i32* [[CAPEC_ADDR1:%.+]], align |
207 | // CHECK: [[ARG1:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR1]], align |
208 | // CHECK: [[CEV:%.+]] = load i32, i32* [[CAPE_ADDR2]], align |
209 | // CHECK-64: [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPEC_ADDR2:%.+]] to i32* |
210 | // CHECK-64: store i32 [[CEV]], i32* [[CONV]], align |
211 | // CHECK-32: store i32 [[CEV]], i32* [[CAPEC_ADDR2:%.+]], align |
212 | // CHECK: [[ARG2:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR2]], align |
213 | // CHECK: [[TEAMS:%.+]] = load i32, i32* [[CAPE_ADDR1]], align |
214 | // CHECK: [[TL:%.+]] = load i32, i32* [[CAPE_ADDR2]], align |
215 | // |
216 | // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 2, {{.*}}, i32 [[TEAMS]], i32 [[TL]]) |
217 | // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 |
218 | // CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] |
219 | // |
220 | // CHECK: [[FAIL]] |
221 | // CHECK: call void [[HVT3:@.+]](i[[SZ]] [[ARG1]], i[[SZ]] [[ARG2]]) |
222 | // CHECK: br label {{%?}}[[END]] |
223 | // CHECK: [[END]] |
224 | // |
225 | // |
226 | // |
227 | // CHECK: [[NV:%.+]] = load i32, i32* [[N_ADDR]], align |
228 | // CHECK: [[ADD:%.+]] = add nsw i32 32, [[NV]] |
229 | // CHECK: store i32 [[ADD]], i32* [[CAPE_ADDR:%.+]], align |
230 | // CHECK: [[CEV:%.+]] = load i32, i32* [[CAPE_ADDR]], align |
231 | // CHECK-64: [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPEC_ADDR:%.+]] to i32* |
232 | // CHECK-64: store i32 [[CEV]], i32* [[CONV]], align |
233 | // CHECK-32: store i32 [[CEV]], i32* [[CAPEC_ADDR:%.+]], align |
234 | // CHECK: [[ARG:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR]], align |
235 | // CHECK: [[TL:%.+]] = load i32, i32* [[CAPE_ADDR]], align |
236 | // |
237 | // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 1, {{.*}}, i32 0, i32 [[TL]]) |
238 | // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 |
239 | // CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] |
240 | // |
241 | // CHECK: [[FAIL]] |
242 | // CHECK: call void [[HVT4:@.+]](i[[SZ]] [[ARG]]) |
243 | // CHECK: br label {{%?}}[[END]] |
244 | // CHECK: [[END]] |
245 | // |
246 | |
247 | |
248 | |
249 | |
250 | |
251 | |
252 | // |
253 | // CHECK: define {{.*}}[[FTEMPLATE]] |
254 | // |
255 | // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 0, {{.*}}, i32 0, i32 20) |
256 | // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 |
257 | // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] |
258 | // |
259 | // CHECK: [[FAIL]] |
260 | // CHECK: call void [[HVT5:@.+]]() |
261 | // CHECK: br label {{%?}}[[END]] |
262 | // |
263 | // CHECK: [[END]] |
264 | // |
265 | // |
266 | // |
267 | // CHECK: store i16 1, i16* [[B:%.+]], align |
268 | // CHECK: [[BV:%.+]] = load i16, i16* [[B]], align |
269 | // CHECK: store i16 [[BV]], i16* [[CAPE_ADDR:%.+]], align |
270 | // CHECK: [[CEV:%.+]] = load i16, i16* [[CAPE_ADDR]], align |
271 | // CHECK: [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPEC_ADDR:%.+]] to i16* |
272 | // CHECK: store i16 [[CEV]], i16* [[CONV]], align |
273 | // CHECK: [[ARG:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR]], align |
274 | // CHECK: [[T:%.+]] = load i16, i16* [[CAPE_ADDR]], align |
275 | // CHECK: [[TEAMS:%.+]] = sext i16 [[T]] to i32 |
276 | // |
277 | // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 3, {{.*}}, i32 [[TEAMS]], i32 1024) |
278 | // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 |
279 | // CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] |
280 | // |
281 | // CHECK: [[FAIL]] |
282 | // CHECK: call void [[HVT6:@.+]](i[[SZ]] {{%.+}}, i[[SZ]] {{%.+}}, i[[SZ]] [[ARG]]) |
283 | // CHECK: br label {{%?}}[[END]] |
284 | // CHECK: [[END]] |
285 | // |
286 | |
287 | |
288 | |
289 | |
290 | |
291 | |
292 | // Check that the offloading functions are emitted and that the parallel function |
293 | // is appropriately guarded. |
294 | |
295 | // CHECK: define internal void [[HVT1]]([[S1]]* {{%.+}}, i[[SZ]] [[PARM1:%.+]], i[[SZ]] [[PARM2:%.+]]) |
296 | // CHECK-DAG: store i[[SZ]] [[PARM2]], i[[SZ]]* [[CAPE_ADDR:%.+]], align |
297 | // CHECK-64: [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPE_ADDR]] to i32* |
298 | // CHECK-64: [[TL:%.+]] = load i32, i32* [[CONV]], align |
299 | // CHECK-32: [[TL:%.+]] = load i32, i32* [[CAPE_ADDR]], align |
300 | // CHECK: call i32 @__kmpc_push_num_teams(%struct.ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 0, i32 [[TL]]) |
301 | // CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* [[DEF_LOC]], i32 2, |
302 | // |
303 | // |
304 | |
305 | |
306 | // CHECK: define internal void [[HVT2]]([[S1]]* {{%.+}}) |
307 | // CHECK: call i32 @__kmpc_push_num_teams(%struct.ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 0, i32 1024) |
308 | // CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* [[DEF_LOC]], i32 1, |
309 | // |
310 | // |
311 | |
312 | |
313 | |
314 | |
315 | |
316 | |
317 | |
318 | |
319 | // CHECK: define internal void [[HVT3]](i[[SZ]] [[PARM1:%.+]], i[[SZ]] [[PARM2:%.+]]) |
320 | // CHECK-DAG: store i[[SZ]] [[PARM1]], i[[SZ]]* [[CAPE_ADDR1:%.+]], align |
321 | // CHECK-DAG: store i[[SZ]] [[PARM2]], i[[SZ]]* [[CAPE_ADDR2:%.+]], align |
322 | // CHECK-64: [[CONV1:%.+]] = bitcast i[[SZ]]* [[CAPE_ADDR1]] to i32* |
323 | // CHECK-64: [[CONV2:%.+]] = bitcast i[[SZ]]* [[CAPE_ADDR2]] to i32* |
324 | // CHECK-64: [[NT:%.+]] = load i32, i32* [[CONV1]], align |
325 | // CHECK-64: [[TL:%.+]] = load i32, i32* [[CONV2]], align |
326 | // CHECK-32: [[NT:%.+]] = load i32, i32* [[CAPE_ADDR1]], align |
327 | // CHECK-32: [[TL:%.+]] = load i32, i32* [[CAPE_ADDR2]], align |
328 | // CHECK: call i32 @__kmpc_push_num_teams(%struct.ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 [[NT]], i32 [[TL]]) |
329 | // CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* [[DEF_LOC]], i32 0, |
330 | // |
331 | // |
332 | // CHECK: define internal void [[HVT4]](i[[SZ]] [[PARM:%.+]]) |
333 | // CHECK-DAG: store i[[SZ]] [[PARM]], i[[SZ]]* [[CAPE_ADDR:%.+]], align |
334 | // CHECK-64: [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPE_ADDR]] to i32* |
335 | // CHECK-64: [[TL:%.+]] = load i32, i32* [[CONV]], align |
336 | // CHECK-32: [[TL:%.+]] = load i32, i32* [[CAPE_ADDR]], align |
337 | // CHECK: call i32 @__kmpc_push_num_teams(%struct.ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 0, i32 [[TL]]) |
338 | // CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* [[DEF_LOC]], i32 0, |
339 | // |
340 | // |
341 | |
342 | |
343 | |
344 | |
345 | |
346 | // CHECK: define internal void [[HVT5]]( |
347 | // CHECK: call i32 @__kmpc_push_num_teams(%struct.ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 0, i32 20) |
348 | // CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* [[DEF_LOC]], i32 0, |
349 | // |
350 | // |
351 | |
352 | |
353 | // CHECK: define internal void [[HVT6]](i[[SZ]] [[PARM1:%.+]], i[[SZ]] [[PARM2:%.+]], i[[SZ]] [[PARM3:%.+]]) |
354 | // CHECK-DAG: store i[[SZ]] [[PARM3]], i[[SZ]]* [[CAPE_ADDR:%.+]], align |
355 | // CHECK: [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPE_ADDR]] to i16* |
356 | // CHECK: [[T:%.+]] = load i16, i16* [[CONV]], align |
357 | // CHECK: [[NT:%.+]] = sext i16 [[T]] to i32 |
358 | // CHECK: call i32 @__kmpc_push_num_teams(%struct.ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 [[NT]], i32 1024) |
359 | // CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* [[DEF_LOC]], i32 2, |
360 | // |
361 | // |
362 | |
363 | |
364 | |
365 | #endif |
366 | |