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 // Test target codegen - host bc file has to be created first. 10 // 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 11 // 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 12 // 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 13 // 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 14 // 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 15 // 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 16 // 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 17 // 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 18 19 // expected-no-diagnostics 20 #ifndef HEADER 21 #define HEADER 22 23 // CHECK-DAG: %ident_t = type { i32, i32, i32, i32, i8* } 24 // CHECK-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00" 25 // 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) } 26 27 // CHECK-DAG: [[S1:%.+]] = type { double } 28 // CHECK-DAG: [[ENTTY:%.+]] = type { i8*, i8*, i[[SZ:32|64]], i32, i32 } 29 // CHECK-DAG: [[DEVTY:%.+]] = type { i8*, i8*, [[ENTTY]]*, [[ENTTY]]* } 30 // CHECK-DAG: [[DSCTY:%.+]] = type { i32, [[DEVTY]]*, [[ENTTY]]*, [[ENTTY]]* } 31 32 // TCHECK: [[ENTTY:%.+]] = type { i8*, i8*, i{{32|64}}, i32, i32 } 33 34 // CHECK-DAG: $[[REGFN:\.omp_offloading\..+]] = comdat 35 36 // We have 6 target regions 37 38 // CHECK-DAG: @{{.*}} = private constant i8 0 39 // CHECK-DAG: @{{.*}} = private constant i8 0 40 // CHECK-DAG: @{{.*}} = private constant i8 0 41 // CHECK-DAG: @{{.*}} = private constant i8 0 42 // CHECK-DAG: @{{.*}} = private constant i8 0 43 // CHECK-DAG: @{{.*}} = private constant i8 0 44 45 // TCHECK: @{{.+}} = constant [[ENTTY]] 46 // TCHECK: @{{.+}} = constant [[ENTTY]] 47 // TCHECK: @{{.+}} = constant [[ENTTY]] 48 // TCHECK: @{{.+}} = constant [[ENTTY]] 49 // TCHECK: @{{.+}} = constant [[ENTTY]] 50 // TCHECK: @{{.+}} = constant [[ENTTY]] 51 52 // Check if offloading descriptor is created. 53 // CHECK: [[ENTBEGIN:@.+]] = external constant [[ENTTY]] 54 // CHECK: [[ENTEND:@.+]] = external constant [[ENTTY]] 55 // CHECK: [[DEVBEGIN:@.+]] = external constant i8 56 // CHECK: [[DEVEND:@.+]] = external constant i8 57 // CHECK: [[IMAGES:@.+]] = internal unnamed_addr constant [1 x [[DEVTY]]] [{{.+}} { i8* [[DEVBEGIN]], i8* [[DEVEND]], [[ENTTY]]* [[ENTBEGIN]], [[ENTTY]]* [[ENTEND]] }], comdat($[[REGFN]]) 58 // 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]]) 59 60 // Check target registration is registered as a Ctor. 61 // 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*) }] 62 63 64 template<typename tx> 65 tx ftemplate(int n) { 66 tx a = 0; 67 68 #pragma omp target teams thread_limit(tx(20)) 69 { 70 } 71 72 short b = 1; 73 #pragma omp target teams num_teams(b) thread_limit(1024) 74 { 75 a += b; 76 } 77 78 return a; 79 } 80 81 static 82 int fstatic(int n) { 83 84 #pragma omp target teams num_teams(n) thread_limit(n*32) 85 { 86 } 87 88 #pragma omp target teams thread_limit(32+n) 89 { 90 } 91 92 return n+1; 93 } 94 95 struct S1 { 96 double a; 97 98 int r1(int n){ 99 int b = 1; 100 101 #pragma omp target teams thread_limit(n-b) 102 { 103 this->a = (double)b + 1.5; 104 } 105 106 #pragma omp target teams thread_limit(1024) 107 { 108 this->a = 2.5; 109 } 110 111 return (int)a; 112 } 113 }; 114 115 // CHECK: define {{.*}}@{{.*}}bar{{.*}} 116 int bar(int n){ 117 int a = 0; 118 119 S1 S; 120 // CHECK: call {{.*}}i32 [[FS1:@.+]]([[S1]]* {{.*}}, i32 {{.*}}) 121 a += S.r1(n); 122 123 // CHECK: call {{.*}}i32 [[FSTATIC:@.+]](i32 {{.*}}) 124 a += fstatic(n); 125 126 // CHECK: call {{.*}}i32 [[FTEMPLATE:@.+]](i32 {{.*}}) 127 a += ftemplate<int>(n); 128 129 return a; 130 } 131 132 133 134 // 135 // CHECK: define {{.*}}[[FS1]]([[S1]]* {{%.+}}, i32 {{[^%]*}}[[PARM:%.+]]) 136 // 137 // CHECK-DAG: store i32 [[PARM]], i32* [[N_ADDR:%.+]], align 138 // CHECK: store i32 1, i32* [[B:%.+]], align 139 // CHECK: [[NV:%.+]] = load i32, i32* [[N_ADDR]], align 140 // CHECK: [[BV:%.+]] = load i32, i32* [[B]], align 141 // CHECK: [[SUB:%.+]] = sub nsw i32 [[NV]], [[BV]] 142 // CHECK: store i32 [[SUB]], i32* [[CAPE_ADDR:%.+]], align 143 // CHECK: [[CEV:%.+]] = load i32, i32* [[CAPE_ADDR]], align 144 // CHECK-64: [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPEC_ADDR:%.+]] to i32* 145 // CHECK-64: store i32 [[CEV]], i32* [[CONV]], align 146 // CHECK-32: store i32 [[CEV]], i32* [[CAPEC_ADDR:%.+]], align 147 // CHECK: [[ARG:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR]], align 148 // CHECK: [[TL:%.+]] = load i32, i32* [[CAPE_ADDR]], align 149 // 150 // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 3, {{.*}}, i32 0, i32 [[TL]]) 151 // CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align 152 // CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align 153 // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 154 // CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] 155 // 156 // CHECK: [[FAIL]] 157 // CHECK: call void [[HVT1:@.+]]([[S1]]* {{%.+}}, i[[SZ]] {{%.+}}, i[[SZ]] [[ARG]]) 158 // CHECK: br label {{%?}}[[END]] 159 // CHECK: [[END]] 160 // 161 // 162 // 163 // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 1, {{.+}}, i32 0, i32 1024) 164 // CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align 165 // CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align 166 // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 167 // CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] 168 // 169 // CHECK: [[FAIL]] 170 // CHECK: call void [[HVT2:@.+]]([[S1]]* {{[^,]+}}) 171 // CHECK: br label {{%?}}[[END]] 172 // CHECK: [[END]] 173 // 174 175 176 177 178 179 180 // 181 // CHECK: define {{.*}}[[FSTATIC]](i32 {{[^%]*}}[[PARM:%.+]]) 182 // 183 // CHECK-DAG: store i32 [[PARM]], i32* [[N_ADDR:%.+]], align 184 // CHECK: [[NV:%.+]] = load i32, i32* [[N_ADDR]], align 185 // CHECK: store i32 [[NV]], i32* [[CAPE_ADDR1:%.+]], align 186 // CHECK: [[NV:%.+]] = load i32, i32* [[N_ADDR]], align 187 // CHECK: [[MUL:%.+]] = mul nsw i32 [[NV]], 32 188 // CHECK: store i32 [[MUL]], i32* [[CAPE_ADDR2:%.+]], align 189 // CHECK: [[CEV:%.+]] = load i32, i32* [[CAPE_ADDR1]], align 190 // CHECK-64: [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPEC_ADDR1:%.+]] to i32* 191 // CHECK-64: store i32 [[CEV]], i32* [[CONV]], align 192 // CHECK-32: store i32 [[CEV]], i32* [[CAPEC_ADDR1:%.+]], align 193 // CHECK: [[ARG1:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR1]], align 194 // CHECK: [[CEV:%.+]] = load i32, i32* [[CAPE_ADDR2]], align 195 // CHECK-64: [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPEC_ADDR2:%.+]] to i32* 196 // CHECK-64: store i32 [[CEV]], i32* [[CONV]], align 197 // CHECK-32: store i32 [[CEV]], i32* [[CAPEC_ADDR2:%.+]], align 198 // CHECK: [[ARG2:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR2]], align 199 // CHECK: [[TEAMS:%.+]] = load i32, i32* [[CAPE_ADDR1]], align 200 // CHECK: [[TL:%.+]] = load i32, i32* [[CAPE_ADDR2]], align 201 // 202 // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 2, {{.*}}, i32 [[TEAMS]], i32 [[TL]]) 203 // CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align 204 // CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align 205 // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 206 // CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] 207 // 208 // CHECK: [[FAIL]] 209 // CHECK: call void [[HVT3:@.+]](i[[SZ]] [[ARG1]], i[[SZ]] [[ARG2]]) 210 // CHECK: br label {{%?}}[[END]] 211 // CHECK: [[END]] 212 // 213 // 214 // 215 // CHECK: [[NV:%.+]] = load i32, i32* [[N_ADDR]], align 216 // CHECK: [[ADD:%.+]] = add nsw i32 32, [[NV]] 217 // CHECK: store i32 [[ADD]], i32* [[CAPE_ADDR:%.+]], align 218 // CHECK: [[CEV:%.+]] = load i32, i32* [[CAPE_ADDR]], align 219 // CHECK-64: [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPEC_ADDR:%.+]] to i32* 220 // CHECK-64: store i32 [[CEV]], i32* [[CONV]], align 221 // CHECK-32: store i32 [[CEV]], i32* [[CAPEC_ADDR:%.+]], align 222 // CHECK: [[ARG:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR]], align 223 // CHECK: [[TL:%.+]] = load i32, i32* [[CAPE_ADDR]], align 224 // 225 // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 1, {{.*}}, i32 0, i32 [[TL]]) 226 // CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align 227 // CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align 228 // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 229 // CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] 230 // 231 // CHECK: [[FAIL]] 232 // CHECK: call void [[HVT4:@.+]](i[[SZ]] [[ARG]]) 233 // CHECK: br label {{%?}}[[END]] 234 // CHECK: [[END]] 235 // 236 237 238 239 240 241 242 // 243 // CHECK: define {{.*}}[[FTEMPLATE]] 244 // 245 // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 0, {{.*}}, i32 0, i32 20) 246 // CHECK-NEXT: store i32 [[RET]], i32* [[RHV:%.+]], align 247 // CHECK-NEXT: [[RET2:%.+]] = load i32, i32* [[RHV]], align 248 // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 249 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] 250 // 251 // CHECK: [[FAIL]] 252 // CHECK: call void [[HVT5:@.+]]() 253 // CHECK: br label {{%?}}[[END]] 254 // 255 // CHECK: [[END]] 256 // 257 // 258 // 259 // CHECK: store i16 1, i16* [[B:%.+]], align 260 // CHECK: [[BV:%.+]] = load i16, i16* [[B]], align 261 // CHECK: store i16 [[BV]], i16* [[CAPE_ADDR:%.+]], align 262 // CHECK: [[CEV:%.+]] = load i16, i16* [[CAPE_ADDR]], align 263 // CHECK: [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPEC_ADDR:%.+]] to i16* 264 // CHECK: store i16 [[CEV]], i16* [[CONV]], align 265 // CHECK: [[ARG:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR]], align 266 // CHECK: [[T:%.+]] = load i16, i16* [[CAPE_ADDR]], align 267 // CHECK: [[TEAMS:%.+]] = sext i16 [[T]] to i32 268 // 269 // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 3, {{.*}}, i32 [[TEAMS]], i32 1024) 270 // CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align 271 // CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align 272 // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 273 // CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] 274 // 275 // CHECK: [[FAIL]] 276 // CHECK: call void [[HVT6:@.+]](i[[SZ]] {{%.+}}, i[[SZ]] {{%.+}}, i[[SZ]] [[ARG]]) 277 // CHECK: br label {{%?}}[[END]] 278 // CHECK: [[END]] 279 // 280 281 282 283 284 285 286 // Check that the offloading functions are emitted and that the parallel function 287 // is appropriately guarded. 288 289 // CHECK: define internal void [[HVT1]]([[S1]]* {{%.+}}, i[[SZ]] [[PARM1:%.+]], i[[SZ]] [[PARM2:%.+]]) 290 // CHECK-DAG: store i[[SZ]] [[PARM2]], i[[SZ]]* [[CAPE_ADDR:%.+]], align 291 // CHECK-64: [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPE_ADDR]] to i32* 292 // CHECK-64: [[TL:%.+]] = load i32, i32* [[CONV]], align 293 // CHECK-32: [[TL:%.+]] = load i32, i32* [[CAPE_ADDR]], align 294 // CHECK: call i32 @__kmpc_push_num_teams(%ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 0, i32 [[TL]]) 295 // CHECK: call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC]], i32 2, 296 // 297 // 298 299 300 // CHECK: define internal void [[HVT2]]([[S1]]* {{%.+}}) 301 // CHECK: call i32 @__kmpc_push_num_teams(%ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 0, i32 1024) 302 // CHECK: call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC]], i32 1, 303 // 304 // 305 306 307 308 309 310 311 312 313 // CHECK: define internal void [[HVT3]](i[[SZ]] [[PARM1:%.+]], i[[SZ]] [[PARM2:%.+]]) 314 // CHECK-DAG: store i[[SZ]] [[PARM1]], i[[SZ]]* [[CAPE_ADDR1:%.+]], align 315 // CHECK-DAG: store i[[SZ]] [[PARM2]], i[[SZ]]* [[CAPE_ADDR2:%.+]], align 316 // CHECK-64: [[CONV1:%.+]] = bitcast i[[SZ]]* [[CAPE_ADDR1]] to i32* 317 // CHECK-64: [[CONV2:%.+]] = bitcast i[[SZ]]* [[CAPE_ADDR2]] to i32* 318 // CHECK-64: [[NT:%.+]] = load i32, i32* [[CONV1]], align 319 // CHECK-64: [[TL:%.+]] = load i32, i32* [[CONV2]], align 320 // CHECK-32: [[NT:%.+]] = load i32, i32* [[CAPE_ADDR1]], align 321 // CHECK-32: [[TL:%.+]] = load i32, i32* [[CAPE_ADDR2]], align 322 // CHECK: call i32 @__kmpc_push_num_teams(%ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 [[NT]], i32 [[TL]]) 323 // CHECK: call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC]], i32 0, 324 // 325 // 326 // CHECK: define internal void [[HVT4]](i[[SZ]] [[PARM:%.+]]) 327 // CHECK-DAG: store i[[SZ]] [[PARM]], i[[SZ]]* [[CAPE_ADDR:%.+]], align 328 // CHECK-64: [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPE_ADDR]] to i32* 329 // CHECK-64: [[TL:%.+]] = load i32, i32* [[CONV]], align 330 // CHECK-32: [[TL:%.+]] = load i32, i32* [[CAPE_ADDR]], align 331 // CHECK: call i32 @__kmpc_push_num_teams(%ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 0, i32 [[TL]]) 332 // CHECK: call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC]], i32 0, 333 // 334 // 335 336 337 338 339 340 // CHECK: define internal void [[HVT5]]( 341 // CHECK: call i32 @__kmpc_push_num_teams(%ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 0, i32 20) 342 // CHECK: call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC]], i32 0, 343 // 344 // 345 346 347 // CHECK: define internal void [[HVT6]](i[[SZ]] [[PARM1:%.+]], i[[SZ]] [[PARM2:%.+]], i[[SZ]] [[PARM3:%.+]]) 348 // CHECK-DAG: store i[[SZ]] [[PARM3]], i[[SZ]]* [[CAPE_ADDR:%.+]], align 349 // CHECK: [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPE_ADDR]] to i16* 350 // CHECK: [[T:%.+]] = load i16, i16* [[CONV]], align 351 // CHECK: [[NT:%.+]] = sext i16 [[T]] to i32 352 // CHECK: call i32 @__kmpc_push_num_teams(%ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 [[NT]], i32 1024) 353 // CHECK: call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC]], i32 2, 354 // 355 // 356 357 358 359 #endif 360