Gitiles
Code Review Sign In
gerrit-public.fairphone.software / toolchain / llvm-project / a8a9153a372afcfe2d18399b079bfdc5fd4da626 / . / clang / test / OpenMP / target_parallel_if_codegen.cpp
blob: 04c052ad6d58b02afe4027db222df1739b23bd6d [file] [log] [blame]
Arpith Chacko Jacobfe4890a2017-01-18 20:40:48 +0000[diff] [blame]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
Alexey Bataeva8a9153a2017-12-29 18:07:07 +0000[diff] [blame^]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
Arpith Chacko Jacobfe4890a2017-01-18 20:40:48 +0000[diff] [blame]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
Alexey Bataeva8a9153a2017-12-29 18:07:07 +0000[diff] [blame^]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
Arpith Chacko Jacobfe4890a2017-01-18 20:40:48 +0000[diff] [blame]37// expected-no-diagnostics
38#ifndef HEADER
39#define HEADER
40
41// CHECK-DAG: %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 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) }
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
George Rokos29d0f002017-05-27 03:03:13 +0000[diff] [blame]52// CHECK-DAG: $[[REGFN:\.omp_offloading\..+]] = comdat
53
Arpith Chacko Jacobfe4890a2017-01-18 20:40:48 +0000[diff] [blame]54// We have 6 target regions
55
56// CHECK-DAG: @{{.*}} = private constant i8 0
57// CHECK-DAG: @{{.*}} = private constant i8 0
58// CHECK-DAG: @{{.*}} = private constant i8 0
59// CHECK-DAG: @{{.*}} = private constant i8 0
60// CHECK-DAG: @{{.*}} = private constant i8 0
61// CHECK-DAG: @{{.*}} = private constant i8 0
62
63// TCHECK: @{{.+}} = constant [[ENTTY]]
64// TCHECK: @{{.+}} = constant [[ENTTY]]
65// TCHECK: @{{.+}} = constant [[ENTTY]]
66// TCHECK: @{{.+}} = constant [[ENTTY]]
67// TCHECK: @{{.+}} = constant [[ENTTY]]
68// TCHECK: @{{.+}} = 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:@.+]] = external constant i8
74// CHECK: [[DEVEND:@.+]] = external constant i8
George Rokos29d0f002017-05-27 03:03:13 +0000[diff] [blame]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]])
Arpith Chacko Jacobfe4890a2017-01-18 20:40:48 +0000[diff] [blame]77
78// Check target registration is registered as a Ctor.
George Rokos29d0f002017-05-27 03:03:13 +0000[diff] [blame]79// CHECK: appending global [1 x { i32, void ()*, i8* }] [{ i32, void ()*, i8* } { i32 0, void ()* bitcast (void (i8*)* @[[REGFN]] to void ()*), i8* bitcast (void (i8*)* @[[REGFN]] to i8*) }]
Arpith Chacko Jacobfe4890a2017-01-18 20:40:48 +0000[diff] [blame]80
81
82template<typename tx>
83tx 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
100static
101int 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
114struct 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{{.*}}
135int 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
Arpith Chacko Jacobfe4890a2017-01-18 20:40:48 +0000[diff] [blame]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//
George Rokos63bc9d62017-11-21 18:25:12 +0000[diff] [blame]166// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 3, {{.*}}, i32 1, i32 0)
Alexey Bataev2a007e02017-10-02 14:20:58 +0000[diff] [blame]167// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
Arpith Chacko Jacobfe4890a2017-01-18 20:40:48 +0000[diff] [blame]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]]
George Rokos63bc9d62017-11-21 18:25:12 +0000[diff] [blame]192// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 2, {{.*}}, i32 1, i32 0)
Alexey Bataev2a007e02017-10-02 14:20:58 +0000[diff] [blame]193// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
194// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
Arpith Chacko Jacobfe4890a2017-01-18 20:40:48 +0000[diff] [blame]195// CHECK: [[FAIL]]
196// CHECK: call void [[HVT2:@.+]]([[S1]]* {{%.+}}, i[[SZ]] [[ARG]])
Alexey Bataev2a007e02017-10-02 14:20:58 +0000[diff] [blame]197// CHECK-NEXT: br label %[[END]]
Arpith Chacko Jacobfe4890a2017-01-18 20:40:48 +0000[diff] [blame]198// CHECK: [[END]]
Alexey Bataev2a007e02017-10-02 14:20:58 +0000[diff] [blame]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]]
Arpith Chacko Jacobfe4890a2017-01-18 20:40:48 +0000[diff] [blame]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]]
George Rokos63bc9d62017-11-21 18:25:12 +0000[diff] [blame]224// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 1, {{.*}}, i32 1, i32 0)
Alexey Bataev2a007e02017-10-02 14:20:58 +0000[diff] [blame]225// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
226// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
Arpith Chacko Jacobfe4890a2017-01-18 20:40:48 +0000[diff] [blame]227// CHECK: [[FAIL]]
228// CHECK: call void [[HVT3:@.+]](i[[SZ]] [[ARG]])
Alexey Bataev2a007e02017-10-02 14:20:58 +0000[diff] [blame]229// CHECK-NEXT: br label %[[END]]
Arpith Chacko Jacobfe4890a2017-01-18 20:40:48 +0000[diff] [blame]230// CHECK: [[END]]
Alexey Bataev2a007e02017-10-02 14:20:58 +0000[diff] [blame]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]]
Arpith Chacko Jacobfe4890a2017-01-18 20:40:48 +0000[diff] [blame]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]]
George Rokos63bc9d62017-11-21 18:25:12 +0000[diff] [blame]245// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 0, {{.*}}, i32 1, i32 0)
Alexey Bataev2a007e02017-10-02 14:20:58 +0000[diff] [blame]246// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
247// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
Arpith Chacko Jacobfe4890a2017-01-18 20:40:48 +0000[diff] [blame]248// CHECK: [[FAIL]]
249// CHECK: call void [[HVT4:@.+]]()
Alexey Bataev2a007e02017-10-02 14:20:58 +0000[diff] [blame]250// CHECK-NEXT: br label %[[END]]
Arpith Chacko Jacobfe4890a2017-01-18 20:40:48 +0000[diff] [blame]251// CHECK: [[END]]
Alexey Bataev2a007e02017-10-02 14:20:58 +0000[diff] [blame]252// CHECK-NEXT: br label %[[IFEND:.+]]
253// CHECK: [[IF_ELSE]]
254// CHECK: call void [[HVT4]]()
255// CHECK-NEXT: br label %[[IFEND]]
256// CHECK: [[IFEND]]
Arpith Chacko Jacobfe4890a2017-01-18 20:40:48 +0000[diff] [blame]257
258
259
260
261
262
263//
264// CHECK: define {{.*}}[[FTEMPLATE]]
265//
George Rokos63bc9d62017-11-21 18:25:12 +0000[diff] [blame]266// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 1, {{.*}}, i32 1, i32 0)
Alexey Bataev2a007e02017-10-02 14:20:58 +0000[diff] [blame]267// CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
Arpith Chacko Jacobfe4890a2017-01-18 20:40:48 +0000[diff] [blame]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//
George Rokos63bc9d62017-11-21 18:25:12 +0000[diff] [blame]278// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 2, {{.*}}, i32 1, i32 0)
Alexey Bataev2a007e02017-10-02 14:20:58 +0000[diff] [blame]279// CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
Arpith Chacko Jacobfe4890a2017-01-18 20:40:48 +0000[diff] [blame]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 (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%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 (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%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 (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%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 (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%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 (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%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