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 // RUN: %clang_cc1 -verify -fopenmp -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
38 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
39 // RUN: %clang_cc1 -fopenmp -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
40 // RUN: %clang_cc1 -verify -fopenmp -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
41 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
42 // RUN: %clang_cc1 -fopenmp -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
43 
44 // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
45 // RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
46 // RUN: %clang_cc1 -fopenmp-simd -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
47 // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
48 // RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
49 // RUN: %clang_cc1 -fopenmp-simd -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
50 // SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
51 
52 // Test target codegen - host bc file has to be created first.
53 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc
54 // RUN: %clang_cc1 -verify -fopenmp -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
55 // RUN: %clang_cc1 -fopenmp -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
56 // RUN: %clang_cc1 -fopenmp -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
57 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc
58 // RUN: %clang_cc1 -verify -fopenmp -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
59 // RUN: %clang_cc1 -fopenmp -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
60 // RUN: %clang_cc1 -fopenmp -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
61 
62 // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc
63 // RUN: %clang_cc1 -verify -fopenmp-simd -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
64 // RUN: %clang_cc1 -fopenmp-simd -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
65 // RUN: %clang_cc1 -fopenmp-simd -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
66 // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc
67 // RUN: %clang_cc1 -verify -fopenmp-simd -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
68 // RUN: %clang_cc1 -fopenmp-simd -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
69 // RUN: %clang_cc1 -fopenmp-simd -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
70 // SIMD-ONLY1-NOT: {{__kmpc|__tgt}}
71 
72 // expected-no-diagnostics
73 #ifndef HEADER
74 #define HEADER
75 
76 // CHECK-DAG: %struct.ident_t = type { i32, i32, i32, i32, i8* }
77 // CHECK-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
78 // CHECK-DAG: [[DEF_LOC:@.+]] = private unnamed_addr constant %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) }
79 
80 // CHECK-DAG: [[S1:%.+]] = type { double }
81 // CHECK-DAG: [[ENTTY:%.+]] = type { i8*, i8*, i[[SZ:32|64]], i32, i32 }
82 
83 // TCHECK: [[ENTTY:%.+]] = type { i8*, i8*, i{{32|64}}, i32, i32 }
84 
85 // We have 6 target regions
86 
87 // CHECK-DAG: @{{.*}} = weak constant i8 0
88 // CHECK-DAG: @{{.*}} = weak constant i8 0
89 // CHECK-DAG: @{{.*}} = weak constant i8 0
90 // CHECK-DAG: @{{.*}} = weak constant i8 0
91 // CHECK-DAG: @{{.*}} = weak constant i8 0
92 // CHECK-DAG: @{{.*}} = weak constant i8 0
93 
94 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
95 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
96 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
97 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
98 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
99 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
100 
101 // Check target registration is registered as a Ctor.
102 // CHECK: appending global [1 x { i32, void ()*, i8* }] [{ i32, void ()*, i8* } { i32 0, void ()* @.omp_offloading.requires_reg, i8* null }]
103 
104 
105 template<typename tx>
106 tx ftemplate(int n) {
107   tx a = 0;
108 
109   #pragma omp target parallel if(parallel: 0)
110   {
111     a += 1;
112   }
113 
114   short b = 1;
115   #pragma omp target parallel if(parallel: 1)
116   {
117     a += b;
118   }
119 
120   return a;
121 }
122 
123 static
124 int fstatic(int n) {
125 
126   #pragma omp target parallel if(n>1)
127   {
128   }
129 
130   #pragma omp target parallel if(target: n-2>2)
131   {
132   }
133 
134   return n+1;
135 }
136 
137 struct S1 {
138   double a;
139 
140   int r1(int n){
141     int b = 1;
142 
143     #pragma omp target parallel if(parallel: n>3)
144     {
145       this->a = (double)b + 1.5;
146     }
147 
148     #pragma omp target parallel if(target: n>4) if(parallel: n>5)
149     {
150       this->a = 2.5;
151     }
152 
153     return (int)a;
154   }
155 };
156 
157 // CHECK: define {{.*}}@{{.*}}bar{{.*}}
158 int bar(int n){
159   int a = 0;
160 
161   S1 S;
162   // CHECK: call {{.*}}i32 [[FS1:@.+]]([[S1]]* {{.*}}, i32 {{.*}})
163   a += S.r1(n);
164 
165   // CHECK: call {{.*}}i32 [[FSTATIC:@.+]](i32 {{.*}})
166   a += fstatic(n);
167 
168   // CHECK: call {{.*}}i32 [[FTEMPLATE:@.+]](i32 {{.*}})
169   a += ftemplate<int>(n);
170 
171   return a;
172 }
173 
174 //
175 // CHECK: define {{.*}}[[FS1]]([[S1]]* {{%.+}}, i32 {{[^%]*}}[[PARM:%.+]])
176 //
177 // CHECK-DAG:   store i32 [[PARM]], i32* [[N_ADDR:%.+]], align
178 // CHECK:       [[NV:%.+]] = load i32, i32* [[N_ADDR]], align
179 // CHECK:       [[CMP:%.+]] = icmp sgt i32 [[NV]], 3
180 // CHECK:       [[FB:%.+]] = zext i1 [[CMP]] to i8
181 // CHECK:       store i8 [[FB]], i8* [[CAPE_ADDR:%.+]], align
182 // CHECK:       [[CAPE:%.+]] = load i8, i8* [[CAPE_ADDR]], align
183 // CHECK:       [[TB:%.+]] = trunc i8 [[CAPE]] to i1
184 // CHECK:       [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPEC_ADDR:%.+]] to i8*
185 // CHECK:       [[FB:%.+]] = zext i1 [[TB]] to i8
186 // CHECK:       store i8 [[FB]], i8* [[CONV]], align
187 // CHECK:       [[ARG:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR]], align
188 //
189 // CHECK-DAG:   [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 4, {{.*}}, i8** null, i32 1, i32 [[NT:%.+]])
190 // CHECK-DAG:   [[NT]] = select i1 %{{.+}}, i32 0, i32 1
191 // CHECK:       [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
192 // CHECK:       br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
193 //
194 // CHECK:       [[FAIL]]
195 // CHECK:       call void [[HVT1:@.+]]([[S1]]* {{%.+}}, i[[SZ]] {{%.+}}, i[[SZ]] [[ARG]])
196 // CHECK:       br label {{%?}}[[END]]
197 // CHECK:       [[END]]
198 //
199 //
200 //
201 // CHECK:       [[NV:%.+]] = load i32, i32* [[N_ADDR]], align
202 // CHECK:       [[CMP:%.+]] = icmp sgt i32 [[NV]], 5
203 // CHECK:       [[FB:%.+]] = zext i1 [[CMP]] to i8
204 // CHECK:       store i8 [[FB]], i8* [[CAPE_ADDR:%.+]], align
205 // CHECK:       [[CAPE:%.+]] = load i8, i8* [[CAPE_ADDR]], align
206 // CHECK:       [[TB:%.+]] = trunc i8 [[CAPE]] to i1
207 // CHECK:       [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPEC_ADDR:%.+]] to i8*
208 // CHECK:       [[FB:%.+]] = zext i1 [[TB]] to i8
209 // CHECK:       store i8 [[FB]], i8* [[CONV]], align
210 // CHECK:       [[ARG:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR]], align
211 // CHECK:       [[NV:%.+]] = load i32, i32* [[N_ADDR]], align
212 // CHECK:       [[CMP:%.+]] = icmp sgt i32 [[NV]], 4
213 // CHECK:       br i1 [[CMP]], label {{%?}}[[IF_THEN:.+]], label {{%?}}[[IF_ELSE:.+]]
214 //
215 // CHECK:       [[IF_THEN]]
216 // CHECK-DAG:   [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 3, {{.*}}, i8** null, i32 1, i32 [[NT:%.+]])
217 // CHECK-DAG:   [[NT]] = select i1 %{{.+}}, i32 0, i32 1
218 // CHECK:       [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
219 // CHECK-NEXT:  br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
220 // CHECK:       [[FAIL]]
221 // CHECK:       call void [[HVT2:@.+]]([[S1]]* {{%.+}}, i[[SZ]] [[ARG]])
222 // CHECK-NEXT:  br label %[[END]]
223 // CHECK:       [[END]]
224 // CHECK-NEXT:  br label %[[IFEND:.+]]
225 // CHECK:       [[IF_ELSE]]
226 // CHECK:       call void [[HVT2]]([[S1]]* {{%.+}}, i[[SZ]] [[ARG]])
227 // CHECK-NEXT:  br label %[[IFEND]]
228 // CHECK:       [[IFEND]]
229 
230 //
231 // CHECK: define {{.*}}[[FSTATIC]](i32 {{[^%]*}}[[PARM:%.+]])
232 //
233 // CHECK-DAG:   store i32 [[PARM]], i32* [[N_ADDR:%.+]], align
234 // CHECK:       [[NV:%.+]] = load i32, i32* [[N_ADDR]], align
235 // CHECK:       [[CMP:%.+]] = icmp sgt i32 [[NV]], 1
236 // CHECK:       [[FB:%.+]] = zext i1 [[CMP]] to i8
237 // CHECK:       store i8 [[FB]], i8* [[CAPE_ADDR:%.+]], align
238 // CHECK:       [[CAPE:%.+]] = load i8, i8* [[CAPE_ADDR]], align
239 // CHECK:       [[TB:%.+]] = trunc i8 [[CAPE]] to i1
240 // CHECK:       [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPEC_ADDR:%.+]] to i8*
241 // CHECK:       [[FB:%.+]] = zext i1 [[TB]] to i8
242 // CHECK:       store i8 [[FB]], i8* [[CONV]], align
243 // CHECK:       [[ARG:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR]], align
244 // CHECK:       [[CAPE2:%.+]] = load i8, i8* [[CAPE_ADDR]], align
245 // CHECK:       [[TB:%.+]] = trunc i8 [[CAPE2]] to i1
246 // CHECK:       br i1 [[TB]], label {{%?}}[[IF_THEN:.+]], label {{%?}}[[IF_ELSE:.+]]
247 //
248 // CHECK:       [[IF_THEN]]
249 // CHECK-DAG:   [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 1, {{.*}}, i8** null, i32 1, i32 [[NT:%.+]])
250 // CHECK-DAG:   [[NT]] = select i1 %{{.+}}, i32 0, i32 1
251 // CHECK:       [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
252 // CHECK-NEXT:  br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
253 // CHECK:       [[FAIL]]
254 // CHECK:       call void [[HVT3:@.+]](i[[SZ]] [[ARG]])
255 // CHECK-NEXT:  br label %[[END]]
256 // CHECK:       [[END]]
257 // CHECK-NEXT:  br label %[[IFEND:.+]]
258 // CHECK:       [[IF_ELSE]]
259 // CHECK:       call void [[HVT3]](i[[SZ]] [[ARG]])
260 // CHECK-NEXT:  br label %[[IFEND]]
261 // CHECK:       [[IFEND]]
262 //
263 //
264 //
265 // CHECK-DAG:   [[NV:%.+]] = load i32, i32* [[N_ADDR]], align
266 // CHECK:       [[SUB:%.+]] = sub nsw i32 [[NV]], 2
267 // CHECK:       [[CMP:%.+]] = icmp sgt i32 [[SUB]], 2
268 // CHECK:       br i1 [[CMP]], label {{%?}}[[IF_THEN:.+]], label {{%?}}[[IF_ELSE:.+]]
269 //
270 // CHECK:       [[IF_THEN]]
271 // CHECK-DAG:   [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 0, {{.*}}, i8** null, i32 1, i32 0)
272 // CHECK:       [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
273 // CHECK-NEXT:  br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
274 // CHECK:       [[FAIL]]
275 // CHECK:       call void [[HVT4:@.+]]()
276 // CHECK-NEXT:  br label %[[END]]
277 // CHECK:       [[END]]
278 // CHECK-NEXT:  br label %[[IFEND:.+]]
279 // CHECK:       [[IF_ELSE]]
280 // CHECK:       call void [[HVT4]]()
281 // CHECK-NEXT:  br label %[[IFEND]]
282 // CHECK:       [[IFEND]]
283 
284 
285 
286 
287 
288 
289 //
290 // CHECK: define {{.*}}[[FTEMPLATE]]
291 //
292 // CHECK-DAG:   [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 1, {{.*}}, i8** null, i32 1, i32 1)
293 // CHECK-NEXT:  [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
294 // CHECK-NEXT:  br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
295 //
296 // CHECK:       [[FAIL]]
297 // CHECK:       call void [[HVT5:@.+]]({{[^,]+}})
298 // CHECK:       br label {{%?}}[[END]]
299 //
300 // CHECK:       [[END]]
301 //
302 //
303 //
304 // CHECK-DAG:   [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 2, {{.*}}, i8** null, i32 1, i32 0)
305 // CHECK-NEXT:  [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
306 // CHECK-NEXT:  br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
307 //
308 // CHECK:       [[FAIL]]
309 // CHECK:       call void [[HVT6:@.+]]({{[^,]+}}, {{[^,]+}})
310 // CHECK:       br label {{%?}}[[END]]
311 // CHECK:       [[END]]
312 
313 
314 
315 
316 
317 
318 // Check that the offloading functions are emitted and that the parallel function
319 // is appropriately guarded.
320 
321 // CHECK:       define internal void [[HVT1]]([[S1]]* {{%.+}}, i[[SZ]] [[PARM1:%.+]], i[[SZ]] [[PARM2:%.+]])
322 // CHECK-DAG:   store i[[SZ]] [[PARM1]], i[[SZ]]* [[B_ADDR:%.+]], align
323 // CHECK-DAG:   store i[[SZ]] [[PARM2]], i[[SZ]]* [[CAPE_ADDR:%.+]], align
324 // CHECK-64:    [[CONVB:%.+]] = bitcast i[[SZ]]* [[B_ADDR]] to i32*
325 // CHECK:       [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPE_ADDR]] to i8*
326 // CHECK-64:    [[BV:%.+]] = load i32, i32* [[CONVB]], align
327 // CHECK-32:    [[BV:%.+]] = load i32, i32* [[B_ADDR]], align
328 // CHECK-64:    [[BC:%.+]] = bitcast i64* [[ARGA:%.+]] to i32*
329 // CHECK-64:    store i32 [[BV]], i32* [[BC]], align
330 // CHECK-64:    [[ARG:%.+]] = load i[[SZ]], i[[SZ]]* [[ARGA]], align
331 // CHECK-32:    store i32 [[BV]], i32* [[ARGA:%.+]], align
332 // CHECK-32:    [[ARG:%.+]] = load i[[SZ]], i[[SZ]]* [[ARGA]], align
333 // CHECK:       [[IFC:%.+]] = load i8, i8* [[CONV]], align
334 // CHECK:       [[TB:%.+]] = trunc i8 [[IFC]] to i1
335 // CHECK:       br i1 [[TB]], label {{%?}}[[IF_THEN:.+]], label {{%?}}[[IF_ELSE:.+]]
336 //
337 // CHECK:       [[IF_THEN]]
338 // 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]])
339 // CHECK:       br label {{%?}}[[END:.+]]
340 //
341 // CHECK:       [[IF_ELSE]]
342 // CHECK:       call void @__kmpc_serialized_parallel(
343 // CHECK:       call void [[OMP_OUTLINED3]](i32* {{%.+}}, i32* {{%.+}}, [[S1]]* {{.+}}, i[[SZ]] [[ARG]])
344 // CHECK:       call void @__kmpc_end_serialized_parallel(
345 // CHECK:       br label {{%?}}[[END]]
346 //
347 // CHECK:       [[END]]
348 //
349 //
350 
351 
352 // CHECK:       define internal void [[HVT2]]([[S1]]* {{%.+}}, i[[SZ]] [[PARM:%.+]])
353 // CHECK-DAG:   store i[[SZ]] [[PARM]], i[[SZ]]* [[CAPE_ADDR:%.+]], align
354 // CHECK:       [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPE_ADDR]] to i8*
355 // CHECK:       [[IFC:%.+]] = load i8, i8* [[CONV]], align
356 // CHECK:       [[TB:%.+]] = trunc i8 [[IFC]] to i1
357 // CHECK:       br i1 [[TB]], label {{%?}}[[IF_THEN:.+]], label {{%?}}[[IF_ELSE:.+]]
358 //
359 // CHECK:       [[IF_THEN]]
360 // 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]]* {{.+}})
361 // CHECK:       br label {{%?}}[[END:.+]]
362 //
363 // CHECK:       [[IF_ELSE]]
364 // CHECK:       call void @__kmpc_serialized_parallel(
365 // CHECK:       call void [[OMP_OUTLINED4]](i32* {{%.+}}, i32* {{%.+}}, [[S1]]* {{.+}})
366 // CHECK:       call void @__kmpc_end_serialized_parallel(
367 // CHECK:       br label {{%?}}[[END]]
368 //
369 // CHECK:       [[END]]
370 //
371 //
372 
373 
374 
375 
376 
377 
378 
379 
380 // CHECK:       define internal void [[HVT3]](i[[SZ]] [[PARM:%.+]])
381 // CHECK-DAG:   store i[[SZ]] [[PARM]], i[[SZ]]* [[CAPE_ADDR:%.+]], align
382 // CHECK:       [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPE_ADDR]] to i8*
383 // CHECK:       [[IFC:%.+]] = load i8, i8* [[CONV]], align
384 // CHECK:       [[TB:%.+]] = trunc i8 [[IFC]] to i1
385 // CHECK:       br i1 [[TB]], label {{%?}}[[IF_THEN:.+]], label {{%?}}[[IF_ELSE:.+]]
386 //
387 // CHECK:       [[IF_THEN]]
388 // 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*, ...)*))
389 // CHECK:       br label {{%?}}[[END:.+]]
390 //
391 // CHECK:       [[IF_ELSE]]
392 // CHECK:       call void @__kmpc_serialized_parallel(
393 // CHECK:       call void [[OMP_OUTLINED1]](i32* {{%.+}}, i32* {{%.+}})
394 // CHECK:       call void @__kmpc_end_serialized_parallel(
395 // CHECK:       br label {{%?}}[[END]]
396 //
397 // CHECK:       [[END]]
398 //
399 //
400 // CHECK:       define internal void [[HVT4]]()
401 // 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*, ...)*))
402 // CHECK-NEXT:  ret
403 //
404 //
405 
406 
407 
408 
409 
410 // CHECK:       define internal void [[HVT5]](
411 // CHECK-NOT:   @__kmpc_fork_call
412 // CHECK:       call void @__kmpc_serialized_parallel(
413 // CHECK:       call void [[OMP_OUTLINED5:@.+]](i32* {{%.+}}, i32* {{%.+}}, i[[SZ]] {{.+}})
414 // CHECK:       call void @__kmpc_end_serialized_parallel(
415 // CHECK:       ret
416 //
417 //
418 
419 
420 // CHECK:       define internal void [[HVT6]](
421 // CHECK-NOT:   call void @__kmpc_serialized_parallel(
422 // CHECK-NOT:   call void [[OMP_OUTLINED5:@.+]](i32* {{%.+}}, i32* {{%.+}}, i[[SZ]] {{.+}})
423 // CHECK-NOT:   call void @__kmpc_end_serialized_parallel(
424 // 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*, ...)*),
425 // CHECK:       ret
426 //
427 //
428 
429 
430 
431 #endif
432