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