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 parallel if(parallel: 0) |
87 | { |
88 | a += 1; |
89 | } |
90 | |
91 | short b = 1; |
92 | #pragma omp target parallel if(parallel: 1) |
93 | { |
94 | a += b; |
95 | } |
96 | |
97 | return a; |
98 | } |
99 | |
100 | static |
101 | int fstatic(int n) { |
102 | |
103 | #pragma omp target parallel if(n>1) |
104 | { |
105 | } |
106 | |
107 | #pragma omp target parallel if(target: n-2>2) |
108 | { |
109 | } |
110 | |
111 | return n+1; |
112 | } |
113 | |
114 | struct S1 { |
115 | double a; |
116 | |
117 | int r1(int n){ |
118 | int b = 1; |
119 | |
120 | #pragma omp target parallel if(parallel: n>3) |
121 | { |
122 | this->a = (double)b + 1.5; |
123 | } |
124 | |
125 | #pragma omp target parallel if(target: n>4) if(parallel: n>5) |
126 | { |
127 | this->a = 2.5; |
128 | } |
129 | |
130 | return (int)a; |
131 | } |
132 | }; |
133 | |
134 | // CHECK: define {{.*}}@{{.*}}bar{{.*}} |
135 | int bar(int n){ |
136 | int a = 0; |
137 | |
138 | S1 S; |
139 | // CHECK: call {{.*}}i32 [[FS1:@.+]]([[S1]]* {{.*}}, i32 {{.*}}) |
140 | a += S.r1(n); |
141 | |
142 | // CHECK: call {{.*}}i32 [[FSTATIC:@.+]](i32 {{.*}}) |
143 | a += fstatic(n); |
144 | |
145 | // CHECK: call {{.*}}i32 [[FTEMPLATE:@.+]](i32 {{.*}}) |
146 | a += ftemplate<int>(n); |
147 | |
148 | return a; |
149 | } |
150 | |
151 | // |
152 | // CHECK: define {{.*}}[[FS1]]([[S1]]* {{%.+}}, i32 {{[^%]*}}[[PARM:%.+]]) |
153 | // |
154 | // CHECK-DAG: store i32 [[PARM]], i32* [[N_ADDR:%.+]], align |
155 | // CHECK: [[NV:%.+]] = load i32, i32* [[N_ADDR]], align |
156 | // CHECK: [[CMP:%.+]] = icmp sgt i32 [[NV]], 3 |
157 | // CHECK: [[FB:%.+]] = zext i1 [[CMP]] to i8 |
158 | // CHECK: store i8 [[FB]], i8* [[CAPE_ADDR:%.+]], align |
159 | // CHECK: [[CAPE:%.+]] = load i8, i8* [[CAPE_ADDR]], align |
160 | // CHECK: [[TB:%.+]] = trunc i8 [[CAPE]] to i1 |
161 | // CHECK: [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPEC_ADDR:%.+]] to i8* |
162 | // CHECK: [[FB:%.+]] = zext i1 [[TB]] to i8 |
163 | // CHECK: store i8 [[FB]], i8* [[CONV]], align |
164 | // CHECK: [[ARG:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR]], align |
165 | // |
166 | // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 4, {{.*}}, i32 1, i32 0) |
167 | // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 |
168 | // CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] |
169 | // |
170 | // CHECK: [[FAIL]] |
171 | // CHECK: call void [[HVT1:@.+]]([[S1]]* {{%.+}}, i[[SZ]] {{%.+}}, i[[SZ]] [[ARG]]) |
172 | // CHECK: br label {{%?}}[[END]] |
173 | // CHECK: [[END]] |
174 | // |
175 | // |
176 | // |
177 | // CHECK: [[NV:%.+]] = load i32, i32* [[N_ADDR]], align |
178 | // CHECK: [[CMP:%.+]] = icmp sgt i32 [[NV]], 5 |
179 | // CHECK: [[FB:%.+]] = zext i1 [[CMP]] to i8 |
180 | // CHECK: store i8 [[FB]], i8* [[CAPE_ADDR:%.+]], align |
181 | // CHECK: [[CAPE:%.+]] = load i8, i8* [[CAPE_ADDR]], align |
182 | // CHECK: [[TB:%.+]] = trunc i8 [[CAPE]] to i1 |
183 | // CHECK: [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPEC_ADDR:%.+]] to i8* |
184 | // CHECK: [[FB:%.+]] = zext i1 [[TB]] to i8 |
185 | // CHECK: store i8 [[FB]], i8* [[CONV]], align |
186 | // CHECK: [[ARG:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR]], align |
187 | // CHECK: [[NV:%.+]] = load i32, i32* [[N_ADDR]], align |
188 | // CHECK: [[CMP:%.+]] = icmp sgt i32 [[NV]], 4 |
189 | // CHECK: br i1 [[CMP]], label {{%?}}[[IF_THEN:.+]], label {{%?}}[[IF_ELSE:.+]] |
190 | // |
191 | // CHECK: [[IF_THEN]] |
192 | // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 3, {{.*}}, i32 1, i32 0) |
193 | // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 |
194 | // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] |
195 | // CHECK: [[FAIL]] |
196 | // CHECK: call void [[HVT2:@.+]]([[S1]]* {{%.+}}, i[[SZ]] [[ARG]]) |
197 | // CHECK-NEXT: br label %[[END]] |
198 | // CHECK: [[END]] |
199 | // CHECK-NEXT: br label %[[IFEND:.+]] |
200 | // CHECK: [[IF_ELSE]] |
201 | // CHECK: call void [[HVT2]]([[S1]]* {{%.+}}, i[[SZ]] [[ARG]]) |
202 | // CHECK-NEXT: br label %[[IFEND]] |
203 | // CHECK: [[IFEND]] |
204 | |
205 | // |
206 | // CHECK: define {{.*}}[[FSTATIC]](i32 {{[^%]*}}[[PARM:%.+]]) |
207 | // |
208 | // CHECK-DAG: store i32 [[PARM]], i32* [[N_ADDR:%.+]], align |
209 | // CHECK: [[NV:%.+]] = load i32, i32* [[N_ADDR]], align |
210 | // CHECK: [[CMP:%.+]] = icmp sgt i32 [[NV]], 1 |
211 | // CHECK: [[FB:%.+]] = zext i1 [[CMP]] to i8 |
212 | // CHECK: store i8 [[FB]], i8* [[CAPE_ADDR:%.+]], align |
213 | // CHECK: [[CAPE:%.+]] = load i8, i8* [[CAPE_ADDR]], align |
214 | // CHECK: [[TB:%.+]] = trunc i8 [[CAPE]] to i1 |
215 | // CHECK: [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPEC_ADDR:%.+]] to i8* |
216 | // CHECK: [[FB:%.+]] = zext i1 [[TB]] to i8 |
217 | // CHECK: store i8 [[FB]], i8* [[CONV]], align |
218 | // CHECK: [[ARG:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR]], align |
219 | // CHECK: [[CAPE2:%.+]] = load i8, i8* [[CAPE_ADDR]], align |
220 | // CHECK: [[TB:%.+]] = trunc i8 [[CAPE2]] to i1 |
221 | // CHECK: br i1 [[TB]], label {{%?}}[[IF_THEN:.+]], label {{%?}}[[IF_ELSE:.+]] |
222 | // |
223 | // CHECK: [[IF_THEN]] |
224 | // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 1, {{.*}}, i32 1, i32 0) |
225 | // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 |
226 | // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] |
227 | // CHECK: [[FAIL]] |
228 | // CHECK: call void [[HVT3:@.+]](i[[SZ]] [[ARG]]) |
229 | // CHECK-NEXT: br label %[[END]] |
230 | // CHECK: [[END]] |
231 | // CHECK-NEXT: br label %[[IFEND:.+]] |
232 | // CHECK: [[IF_ELSE]] |
233 | // CHECK: call void [[HVT3]](i[[SZ]] [[ARG]]) |
234 | // CHECK-NEXT: br label %[[IFEND]] |
235 | // CHECK: [[IFEND]] |
236 | // |
237 | // |
238 | // |
239 | // CHECK-DAG: [[NV:%.+]] = load i32, i32* [[N_ADDR]], align |
240 | // CHECK: [[SUB:%.+]] = sub nsw i32 [[NV]], 2 |
241 | // CHECK: [[CMP:%.+]] = icmp sgt i32 [[SUB]], 2 |
242 | // CHECK: br i1 [[CMP]], label {{%?}}[[IF_THEN:.+]], label {{%?}}[[IF_ELSE:.+]] |
243 | // |
244 | // CHECK: [[IF_THEN]] |
245 | // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 0, {{.*}}, i32 1, i32 0) |
246 | // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 |
247 | // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] |
248 | // CHECK: [[FAIL]] |
249 | // CHECK: call void [[HVT4:@.+]]() |
250 | // CHECK-NEXT: br label %[[END]] |
251 | // CHECK: [[END]] |
252 | // CHECK-NEXT: br label %[[IFEND:.+]] |
253 | // CHECK: [[IF_ELSE]] |
254 | // CHECK: call void [[HVT4]]() |
255 | // CHECK-NEXT: br label %[[IFEND]] |
256 | // CHECK: [[IFEND]] |
257 | |
258 | |
259 | |
260 | |
261 | |
262 | |
263 | // |
264 | // CHECK: define {{.*}}[[FTEMPLATE]] |
265 | // |
266 | // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 1, {{.*}}, i32 1, i32 0) |
267 | // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 |
268 | // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] |
269 | // |
270 | // CHECK: [[FAIL]] |
271 | // CHECK: call void [[HVT5:@.+]]({{[^,]+}}) |
272 | // CHECK: br label {{%?}}[[END]] |
273 | // |
274 | // CHECK: [[END]] |
275 | // |
276 | // |
277 | // |
278 | // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 2, {{.*}}, i32 1, i32 0) |
279 | // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 |
280 | // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] |
281 | // |
282 | // CHECK: [[FAIL]] |
283 | // CHECK: call void [[HVT6:@.+]]({{[^,]+}}, {{[^,]+}}) |
284 | // CHECK: br label {{%?}}[[END]] |
285 | // CHECK: [[END]] |
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]] [[PARM1]], i[[SZ]]* [[B_ADDR:%.+]], align |
297 | // CHECK-DAG: store i[[SZ]] [[PARM2]], i[[SZ]]* [[CAPE_ADDR:%.+]], align |
298 | // CHECK-64: [[CONVB:%.+]] = bitcast i[[SZ]]* [[B_ADDR]] to i32* |
299 | // CHECK: [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPE_ADDR]] to i8* |
300 | // CHECK-64: [[BV:%.+]] = load i32, i32* [[CONVB]], align |
301 | // CHECK-32: [[BV:%.+]] = load i32, i32* [[B_ADDR]], align |
302 | // CHECK-64: [[BC:%.+]] = bitcast i64* [[ARGA:%.+]] to i32* |
303 | // CHECK-64: store i32 [[BV]], i32* [[BC]], align |
304 | // CHECK-64: [[ARG:%.+]] = load i[[SZ]], i[[SZ]]* [[ARGA]], align |
305 | // CHECK-32: store i32 [[BV]], i32* [[ARGA:%.+]], align |
306 | // CHECK-32: [[ARG:%.+]] = load i[[SZ]], i[[SZ]]* [[ARGA]], align |
307 | // CHECK: [[IFC:%.+]] = load i8, i8* [[CONV]], align |
308 | // CHECK: [[TB:%.+]] = trunc i8 [[IFC]] to i1 |
309 | // CHECK: br i1 [[TB]], label {{%?}}[[IF_THEN:.+]], label {{%?}}[[IF_ELSE:.+]] |
310 | // |
311 | // CHECK: [[IF_THEN]] |
312 | // CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* [[DEF_LOC]], i32 2, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, [[S1]]*, i[[SZ]])* [[OMP_OUTLINED3:@.+]] to void (i32*, i32*, ...)*), [[S1]]* {{.+}}, i[[SZ]] [[ARG]]) |
313 | // CHECK: br label {{%?}}[[END:.+]] |
314 | // |
315 | // CHECK: [[IF_ELSE]] |
316 | // CHECK: call void @__kmpc_serialized_parallel( |
317 | // CHECK: call void [[OMP_OUTLINED3]](i32* {{%.+}}, i32* {{%.+}}, [[S1]]* {{.+}}, i[[SZ]] [[ARG]]) |
318 | // CHECK: call void @__kmpc_end_serialized_parallel( |
319 | // CHECK: br label {{%?}}[[END]] |
320 | // |
321 | // CHECK: [[END]] |
322 | // |
323 | // |
324 | |
325 | |
326 | // CHECK: define internal void [[HVT2]]([[S1]]* {{%.+}}, i[[SZ]] [[PARM:%.+]]) |
327 | // CHECK-DAG: store i[[SZ]] [[PARM]], i[[SZ]]* [[CAPE_ADDR:%.+]], align |
328 | // CHECK: [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPE_ADDR]] to i8* |
329 | // CHECK: [[IFC:%.+]] = load i8, i8* [[CONV]], align |
330 | // CHECK: [[TB:%.+]] = trunc i8 [[IFC]] to i1 |
331 | // CHECK: br i1 [[TB]], label {{%?}}[[IF_THEN:.+]], label {{%?}}[[IF_ELSE:.+]] |
332 | // |
333 | // CHECK: [[IF_THEN]] |
334 | // CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* [[DEF_LOC]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, [[S1]]*)* [[OMP_OUTLINED4:@.+]] to void (i32*, i32*, ...)*), [[S1]]* {{.+}}) |
335 | // CHECK: br label {{%?}}[[END:.+]] |
336 | // |
337 | // CHECK: [[IF_ELSE]] |
338 | // CHECK: call void @__kmpc_serialized_parallel( |
339 | // CHECK: call void [[OMP_OUTLINED4]](i32* {{%.+}}, i32* {{%.+}}, [[S1]]* {{.+}}) |
340 | // CHECK: call void @__kmpc_end_serialized_parallel( |
341 | // CHECK: br label {{%?}}[[END]] |
342 | // |
343 | // CHECK: [[END]] |
344 | // |
345 | // |
346 | |
347 | |
348 | |
349 | |
350 | |
351 | |
352 | |
353 | |
354 | // CHECK: define internal void [[HVT3]](i[[SZ]] [[PARM:%.+]]) |
355 | // CHECK-DAG: store i[[SZ]] [[PARM]], i[[SZ]]* [[CAPE_ADDR:%.+]], align |
356 | // CHECK: [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPE_ADDR]] to i8* |
357 | // CHECK: [[IFC:%.+]] = load i8, i8* [[CONV]], align |
358 | // CHECK: [[TB:%.+]] = trunc i8 [[IFC]] to i1 |
359 | // CHECK: br i1 [[TB]], label {{%?}}[[IF_THEN:.+]], label {{%?}}[[IF_ELSE:.+]] |
360 | // |
361 | // CHECK: [[IF_THEN]] |
362 | // CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* [[DEF_LOC]], i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* [[OMP_OUTLINED1:@.+]] to void (i32*, i32*, ...)*)) |
363 | // CHECK: br label {{%?}}[[END:.+]] |
364 | // |
365 | // CHECK: [[IF_ELSE]] |
366 | // CHECK: call void @__kmpc_serialized_parallel( |
367 | // CHECK: call void [[OMP_OUTLINED1]](i32* {{%.+}}, i32* {{%.+}}) |
368 | // CHECK: call void @__kmpc_end_serialized_parallel( |
369 | // CHECK: br label {{%?}}[[END]] |
370 | // |
371 | // CHECK: [[END]] |
372 | // |
373 | // |
374 | // CHECK: define internal void [[HVT4]]() |
375 | // CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* [[DEF_LOC]], i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* [[OMP_OUTLINED2:@.+]] to void (i32*, i32*, ...)*)) |
376 | // CHECK-NEXT: ret |
377 | // |
378 | // |
379 | |
380 | |
381 | |
382 | |
383 | |
384 | // CHECK: define internal void [[HVT5]]( |
385 | // CHECK-NOT: @__kmpc_fork_call |
386 | // CHECK: call void @__kmpc_serialized_parallel( |
387 | // CHECK: call void [[OMP_OUTLINED5:@.+]](i32* {{%.+}}, i32* {{%.+}}, i[[SZ]] {{.+}}) |
388 | // CHECK: call void @__kmpc_end_serialized_parallel( |
389 | // CHECK: ret |
390 | // |
391 | // |
392 | |
393 | |
394 | // CHECK: define internal void [[HVT6]]( |
395 | // CHECK-NOT: call void @__kmpc_serialized_parallel( |
396 | // CHECK-NOT: call void [[OMP_OUTLINED5:@.+]](i32* {{%.+}}, i32* {{%.+}}, i[[SZ]] {{.+}}) |
397 | // CHECK-NOT: call void @__kmpc_end_serialized_parallel( |
398 | // CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* [[DEF_LOC]], i32 2, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i[[SZ]], i[[SZ]])* [[OMP_OUTLINED5:@.+]] to void (i32*, i32*, ...)*), |
399 | // CHECK: ret |
400 | // |
401 | // |
402 | |
403 | |
404 | |
405 | #endif |
406 | |