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 // Test host codegen.
38 // 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
39 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
40 // 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
41 // 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
42 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
43 // 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
44 
45 // 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
46 // 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
47 // 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
48 // 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
49 // 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
50 // 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
51 // SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
52 
53 // Test target codegen - host bc file has to be created first.
54 // 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
55 // 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
56 // 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
57 // 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
58 // 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
59 // 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
60 // 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
61 // 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
62 
63 // 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
64 // 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
65 // 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
66 // 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
67 // 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
68 // 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
69 // 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
70 // 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
71 // SIMD-ONLY1-NOT: {{__kmpc|__tgt}}
72 
73 // expected-no-diagnostics
74 #ifndef HEADER
75 #define HEADER
76 
77 // CHECK-DAG: %struct.ident_t = type { i32, i32, i32, i32, i8* }
78 // CHECK-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
79 // 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) }
80 
81 // CHECK-DAG: [[S1:%.+]] = type { double }
82 // CHECK-DAG: [[ENTTY:%.+]] = type { i8*, i8*, i[[SZ:32|64]], i32, i32 }
83 
84 // TCHECK: [[ENTTY:%.+]] = type { i8*, i8*, i{{32|64}}, i32, i32 }
85 
86 // We have 6 target regions
87 
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 // CHECK-DAG: @{{.*}} = weak constant i8 0
94 
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 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
101 
102 // Check target registration is registered as a Ctor.
103 // CHECK: appending global [1 x { i32, void ()*, i8* }] [{ i32, void ()*, i8* } { i32 0, void ()* @.omp_offloading.requires_reg, i8* null }]
104 
105 
106 template<typename tx>
107 tx ftemplate(int n) {
108   tx a = 0;
109 
110   #pragma omp target teams thread_limit(tx(20))
111   {
112   }
113 
114   short b = 1;
115   #pragma omp target teams num_teams(b) thread_limit(1024)
116   {
117     a += b;
118   }
119 
120   return a;
121 }
122 
123 static
124 int fstatic(int n) {
125 
126   #pragma omp target teams num_teams(n) thread_limit(n*32)
127   {
128   }
129 
130   #pragma omp target teams thread_limit(32+n)
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 teams thread_limit(n-b)
144     {
145       this->a = (double)b + 1.5;
146     }
147 
148     #pragma omp target teams thread_limit(1024)
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 
176 //
177 // CHECK: define {{.*}}[[FS1]]([[S1]]* {{%.+}}, i32 {{[^%]*}}[[PARM:%.+]])
178 //
179 // CHECK-DAG:   store i32 [[PARM]], i32* [[N_ADDR:%.+]], align
180 // CHECK:       store i32 1, i32* [[B:%.+]], align
181 // CHECK:       [[NV:%.+]] = load i32, i32* [[N_ADDR]], align
182 // CHECK:       [[BV:%.+]] = load i32, i32* [[B]], align
183 // CHECK:       [[SUB:%.+]] = sub nsw i32 [[NV]], [[BV]]
184 // CHECK:       store i32 [[SUB]], i32* [[CAPE_ADDR:%.+]], align
185 // CHECK:       [[CEV:%.+]] = load i32, i32* [[CAPE_ADDR]], align
186 // CHECK-64:    [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPEC_ADDR:%.+]] to i32*
187 // CHECK-64:    store i32 [[CEV]], i32* [[CONV]], align
188 // CHECK-32:    store i32 [[CEV]], i32* [[CAPEC_ADDR:%.+]], align
189 // CHECK:       [[ARG:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR]], align
190 // CHECK:       [[TL:%.+]] = load i32, i32* [[CAPE_ADDR]], align
191 //
192 // CHECK-DAG:   [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 4, {{.*}}, i8** null, i32 0, i32 [[TL]])
193 // CHECK:       [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
194 // CHECK:       br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
195 //
196 // CHECK:       [[FAIL]]
197 // CHECK:       call void [[HVT1:@.+]]([[S1]]* {{%.+}}, i[[SZ]] {{%.+}}, i[[SZ]] [[ARG]])
198 // CHECK:       br label {{%?}}[[END]]
199 // CHECK:       [[END]]
200 //
201 //
202 //
203 // CHECK-DAG:   [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 2, {{.+}}, i8** null, i32 0, i32 1024)
204 // CHECK:       [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
205 // CHECK:       br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
206 //
207 // CHECK:       [[FAIL]]
208 // CHECK:       call void [[HVT2:@.+]]([[S1]]* {{[^,]+}})
209 // CHECK:       br label {{%?}}[[END]]
210 // CHECK:       [[END]]
211 //
212 
213 
214 
215 
216 
217 
218 //
219 // CHECK: define {{.*}}[[FSTATIC]](i32 {{[^%]*}}[[PARM:%.+]])
220 //
221 // CHECK-DAG:   store i32 [[PARM]], i32* [[N_ADDR:%.+]], align
222 // CHECK:       [[NV:%.+]] = load i32, i32* [[N_ADDR]], align
223 // CHECK:       store i32 [[NV]], i32* [[CAPE_ADDR1:%.+]], align
224 // CHECK:       [[NV:%.+]] = load i32, i32* [[N_ADDR]], align
225 // CHECK:       [[MUL:%.+]] = mul nsw i32 [[NV]], 32
226 // CHECK:       store i32 [[MUL]], i32* [[CAPE_ADDR2:%.+]], align
227 // CHECK:       [[CEV:%.+]] = load i32, i32* [[CAPE_ADDR1]], align
228 // CHECK-64:    [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPEC_ADDR1:%.+]] to i32*
229 // CHECK-64:    store i32 [[CEV]], i32* [[CONV]], align
230 // CHECK-32:    store i32 [[CEV]], i32* [[CAPEC_ADDR1:%.+]], align
231 // CHECK:       [[ARG1:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR1]], align
232 // CHECK:       [[CEV:%.+]] = load i32, i32* [[CAPE_ADDR2]], align
233 // CHECK-64:    [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPEC_ADDR2:%.+]] to i32*
234 // CHECK-64:    store i32 [[CEV]], i32* [[CONV]], align
235 // CHECK-32:    store i32 [[CEV]], i32* [[CAPEC_ADDR2:%.+]], align
236 // CHECK:       [[ARG2:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR2]], align
237 // CHECK:       [[TEAMS:%.+]] = load i32, i32* [[CAPE_ADDR1]], align
238 // CHECK:       [[TL:%.+]] = load i32, i32* [[CAPE_ADDR2]], align
239 //
240 // CHECK-DAG:   [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 2, {{.*}}, i8** null, i32 [[TEAMS]], i32 [[TL]])
241 // CHECK:       [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
242 // CHECK:       br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
243 //
244 // CHECK:       [[FAIL]]
245 // CHECK:       call void [[HVT3:@.+]](i[[SZ]] [[ARG1]], i[[SZ]] [[ARG2]])
246 // CHECK:       br label {{%?}}[[END]]
247 // CHECK:       [[END]]
248 //
249 //
250 //
251 // CHECK:       [[NV:%.+]] = load i32, i32* [[N_ADDR]], align
252 // CHECK:       [[ADD:%.+]] = add nsw i32 32, [[NV]]
253 // CHECK:       store i32 [[ADD]], i32* [[CAPE_ADDR:%.+]], align
254 // CHECK:       [[CEV:%.+]] = load i32, i32* [[CAPE_ADDR]], align
255 // CHECK-64:    [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPEC_ADDR:%.+]] to i32*
256 // CHECK-64:    store i32 [[CEV]], i32* [[CONV]], align
257 // CHECK-32:    store i32 [[CEV]], i32* [[CAPEC_ADDR:%.+]], align
258 // CHECK:       [[ARG:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR]], align
259 // CHECK:       [[TL:%.+]] = load i32, i32* [[CAPE_ADDR]], align
260 //
261 // CHECK-DAG:   [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 1, {{.*}}, i8** null, i32 0, i32 [[TL]])
262 // CHECK:       [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
263 // CHECK:       br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
264 //
265 // CHECK:       [[FAIL]]
266 // CHECK:       call void [[HVT4:@.+]](i[[SZ]] [[ARG]])
267 // CHECK:       br label {{%?}}[[END]]
268 // CHECK:       [[END]]
269 //
270 
271 
272 
273 
274 
275 
276 //
277 // CHECK: define {{.*}}[[FTEMPLATE]]
278 //
279 // CHECK-DAG:   [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 0, {{.*}}, i8** null, i32 0, i32 20)
280 // CHECK-NEXT:  [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
281 // CHECK-NEXT:  br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
282 //
283 // CHECK:       [[FAIL]]
284 // CHECK:       call void [[HVT5:@.+]]()
285 // CHECK:       br label {{%?}}[[END]]
286 //
287 // CHECK:       [[END]]
288 //
289 //
290 //
291 // CHECK:       store i16 1, i16* [[B:%.+]], align
292 // CHECK:       [[BV:%.+]] = load i16, i16* [[B]], align
293 // CHECK:       store i16 [[BV]], i16* [[CAPE_ADDR:%.+]], align
294 // CHECK:       [[CEV:%.+]] = load i16, i16* [[CAPE_ADDR]], align
295 // CHECK:       [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPEC_ADDR:%.+]] to i16*
296 // CHECK:       store i16 [[CEV]], i16* [[CONV]], align
297 // CHECK:       [[ARG:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR]], align
298 // CHECK:       [[T:%.+]] = load i16, i16* [[CAPE_ADDR]], align
299 // CHECK:       [[TEAMS:%.+]] = sext i16 [[T]] to i32
300 //
301 // CHECK-DAG:   [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 3, {{.*}}, i8** null, i32 [[TEAMS]], i32 1024)
302 // CHECK:       [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
303 // CHECK:       br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
304 //
305 // CHECK:       [[FAIL]]
306 // CHECK:       call void [[HVT6:@.+]](i[[SZ]] {{%.+}}, i[[SZ]] {{%.+}}, i[[SZ]] [[ARG]])
307 // CHECK:       br label {{%?}}[[END]]
308 // CHECK:       [[END]]
309 //
310 
311 
312 
313 
314 
315 
316 // Check that the offloading functions are emitted and that the parallel function
317 // is appropriately guarded.
318 
319 // CHECK:       define internal void [[HVT1]]([[S1]]* {{%.+}}, i[[SZ]] [[PARM1:%.+]], i[[SZ]] [[PARM2:%.+]])
320 // CHECK-DAG:   store i[[SZ]] [[PARM2]], i[[SZ]]* [[CAPE_ADDR:%.+]], align
321 // CHECK-64:    [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPE_ADDR]] to i32*
322 // CHECK-64:    [[TL:%.+]] = load i32, i32* [[CONV]], align
323 // CHECK-32:    [[TL:%.+]] = load i32, i32* [[CAPE_ADDR]], align
324 // CHECK:       call void @__kmpc_push_num_teams(%struct.ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 0, i32 [[TL]])
325 // CHECK:       call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* [[DEF_LOC]], i32 2,
326 //
327 //
328 
329 // CHECK:       define internal void [[HVT2]]([[S1]]* {{%.+}})
330 // CHECK:       call void @__kmpc_push_num_teams(%struct.ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 0, i32 1024)
331 // CHECK:       call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* [[DEF_LOC]], i32 1,
332 //
333 //
334 
335 // CHECK:       define internal void [[HVT3]](i[[SZ]] [[PARM1:%.+]], i[[SZ]] [[PARM2:%.+]])
336 // CHECK-DAG:   store i[[SZ]] [[PARM1]], i[[SZ]]* [[CAPE_ADDR1:%.+]], align
337 // CHECK-DAG:   store i[[SZ]] [[PARM2]], i[[SZ]]* [[CAPE_ADDR2:%.+]], align
338 // CHECK-64:    [[CONV1:%.+]] = bitcast i[[SZ]]* [[CAPE_ADDR1]] to i32*
339 // CHECK-64:    [[CONV2:%.+]] = bitcast i[[SZ]]* [[CAPE_ADDR2]] to i32*
340 // CHECK-64:    [[NT:%.+]] = load i32, i32* [[CONV1]], align
341 // CHECK-64:    [[TL:%.+]] = load i32, i32* [[CONV2]], align
342 // CHECK-32:    [[NT:%.+]] = load i32, i32* [[CAPE_ADDR1]], align
343 // CHECK-32:    [[TL:%.+]] = load i32, i32* [[CAPE_ADDR2]], align
344 // CHECK:       call void @__kmpc_push_num_teams(%struct.ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 [[NT]], i32 [[TL]])
345 // CHECK:       call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* [[DEF_LOC]], i32 0,
346 //
347 //
348 // CHECK:       define internal void [[HVT4]](i[[SZ]] [[PARM:%.+]])
349 // CHECK-DAG:   store i[[SZ]] [[PARM]], i[[SZ]]* [[CAPE_ADDR:%.+]], align
350 // CHECK-64:    [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPE_ADDR]] to i32*
351 // CHECK-64:    [[TL:%.+]] = load i32, i32* [[CONV]], align
352 // CHECK-32:    [[TL:%.+]] = load i32, i32* [[CAPE_ADDR]], align
353 // CHECK:       call void @__kmpc_push_num_teams(%struct.ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 0, i32 [[TL]])
354 // CHECK:       call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* [[DEF_LOC]], i32 0,
355 //
356 //
357 
358 // CHECK:       define internal void [[HVT5]](
359 // CHECK:       call void @__kmpc_push_num_teams(%struct.ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 0, i32 20)
360 // CHECK:       call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* [[DEF_LOC]], i32 0,
361 //
362 //
363 
364 // CHECK:       define internal void [[HVT6]](i[[SZ]] [[PARM1:%.+]], i[[SZ]] [[PARM2:%.+]], i[[SZ]] [[PARM3:%.+]])
365 // CHECK-DAG:   store i[[SZ]] [[PARM3]], i[[SZ]]* [[CAPE_ADDR:%.+]], align
366 // CHECK:       [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPE_ADDR]] to i16*
367 // CHECK:       [[T:%.+]] = load i16, i16* [[CONV]], align
368 // CHECK:       [[NT:%.+]] = sext i16 [[T]] to i32
369 // CHECK:       call void @__kmpc_push_num_teams(%struct.ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 [[NT]], i32 1024)
370 // CHECK:       call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* [[DEF_LOC]], i32 2,
371 //
372 //
373 
374 #endif
375