From: Alexey Bataev Date: Mon, 2 Oct 2017 14:20:58 +0000 (+0000) Subject: [OPENMP] Simplify codegen for non-offloading code. X-Git-Tag: llvmorg-6.0.0-rc1~6757 X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=2a007e05a09986d96dabe1d38a5e7947b20006a2;p=platform%2Fupstream%2Fllvm.git [OPENMP] Simplify codegen for non-offloading code. Simplified and generalized codegen for non-offloading part that works if offloading is failed or condition of the `if` clause is `false`. llvm-svn: 314670 --- diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index fb2ce6c..79675bf 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -6865,8 +6865,6 @@ void CGOpenMPRuntime::emitTargetCall(CodeGenFunction &CGF, assert(OutlinedFn && "Invalid outlined function!"); - auto &Ctx = CGF.getContext(); - // Fill up the arrays with all the captured variables. MappableExprsHandler::MapValuesArrayTy KernelArgs; MappableExprsHandler::MapBaseValuesArrayTy BasePointers; @@ -6931,19 +6929,10 @@ void CGOpenMPRuntime::emitTargetCall(CodeGenFunction &CGF, MapTypes.append(CurMapTypes.begin(), CurMapTypes.end()); } - // Keep track on whether the host function has to be executed. - auto OffloadErrorQType = - Ctx.getIntTypeForBitwidth(/*DestWidth=*/32, /*Signed=*/true); - auto OffloadError = CGF.MakeAddrLValue( - CGF.CreateMemTemp(OffloadErrorQType, ".run_host_version"), - OffloadErrorQType); - CGF.EmitStoreOfScalar(llvm::Constant::getNullValue(CGM.Int32Ty), - OffloadError); - // Fill up the pointer arrays and transfer execution to the device. - auto &&ThenGen = [&BasePointers, &Pointers, &Sizes, &MapTypes, Device, - OutlinedFnID, OffloadError, - &D](CodeGenFunction &CGF, PrePostActionTy &) { + auto &&ThenGen = [this, &BasePointers, &Pointers, &Sizes, &MapTypes, Device, + OutlinedFn, OutlinedFnID, &D, + &KernelArgs](CodeGenFunction &CGF, PrePostActionTy &) { auto &RT = CGF.CGM.getOpenMPRuntime(); // Emit the offloading arrays. TargetDataInfo Info; @@ -7034,13 +7023,26 @@ void CGOpenMPRuntime::emitTargetCall(CodeGenFunction &CGF, OffloadingArgs); } - CGF.EmitStoreOfScalar(Return, OffloadError); + // Check the error code and execute the host version if required. + llvm::BasicBlock *OffloadFailedBlock = + CGF.createBasicBlock("omp_offload.failed"); + llvm::BasicBlock *OffloadContBlock = + CGF.createBasicBlock("omp_offload.cont"); + llvm::Value *Failed = CGF.Builder.CreateIsNotNull(Return); + CGF.Builder.CreateCondBr(Failed, OffloadFailedBlock, OffloadContBlock); + + CGF.EmitBlock(OffloadFailedBlock); + emitOutlinedFunctionCall(CGF, D.getLocStart(), OutlinedFn, KernelArgs); + CGF.EmitBranch(OffloadContBlock); + + CGF.EmitBlock(OffloadContBlock, /*IsFinished=*/true); }; // Notify that the host version must be executed. - auto &&ElseGen = [OffloadError](CodeGenFunction &CGF, PrePostActionTy &) { - CGF.EmitStoreOfScalar(llvm::ConstantInt::get(CGF.Int32Ty, /*V=*/-1u), - OffloadError); + auto &&ElseGen = [this, &D, OutlinedFn, &KernelArgs](CodeGenFunction &CGF, + PrePostActionTy &) { + emitOutlinedFunctionCall(CGF, D.getLocStart(), OutlinedFn, + KernelArgs); }; // If we have a target function ID it means that we need to support @@ -7058,19 +7060,6 @@ void CGOpenMPRuntime::emitTargetCall(CodeGenFunction &CGF, RegionCodeGenTy ElseRCG(ElseGen); ElseRCG(CGF); } - - // Check the error code and execute the host version if required. - auto OffloadFailedBlock = CGF.createBasicBlock("omp_offload.failed"); - auto OffloadContBlock = CGF.createBasicBlock("omp_offload.cont"); - auto OffloadErrorVal = CGF.EmitLoadOfScalar(OffloadError, SourceLocation()); - auto Failed = CGF.Builder.CreateIsNotNull(OffloadErrorVal); - CGF.Builder.CreateCondBr(Failed, OffloadFailedBlock, OffloadContBlock); - - CGF.EmitBlock(OffloadFailedBlock); - emitOutlinedFunctionCall(CGF, D.getLocStart(), OutlinedFn, KernelArgs); - CGF.EmitBranch(OffloadContBlock); - - CGF.EmitBlock(OffloadContBlock, /*IsFinished=*/true); } void CGOpenMPRuntime::scanForTargetRegionsFunctions(const Stmt *S, diff --git a/clang/test/OpenMP/target_codegen.cpp b/clang/test/OpenMP/target_codegen.cpp index 6ed4e59..153f19d 100644 --- a/clang/test/OpenMP/target_codegen.cpp +++ b/clang/test/OpenMP/target_codegen.cpp @@ -93,9 +93,7 @@ int foo(int n) { TT d; // CHECK: [[RET:%.+]] = call i32 @__tgt_target(i32 -1, i8* @{{[^,]+}}, i32 0, i8** null, i8** null, i[[SZ]]* null, i32* null) - // CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align 4 - // CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align 4 - // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 + // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]] // CHECK: [[FAIL]] // CHECK: call void [[HVT0:@.+]]() @@ -105,10 +103,6 @@ int foo(int n) { { } - // CHECK: store i32 0, i32* [[RHV:%.+]], align 4 - // CHECK: store i32 -1, i32* [[RHV]], align 4 - // CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align 4 - // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 // CHECK: call void [[HVT1:@.+]](i[[SZ]] {{[^,]+}}) #pragma omp target if(0) firstprivate(global) { @@ -125,9 +119,7 @@ int foo(int n) { // CHECK-DAG: store i[[SZ]] [[BP0:%[^,]+]], i[[SZ]]* [[CBPADDR0]] // CHECK-DAG: store i[[SZ]] [[P0:%[^,]+]], i[[SZ]]* [[CPADDR0]] - // CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align 4 - // CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align 4 - // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 + // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]] // CHECK: [[FAIL]] // CHECK: call void [[HVT2:@.+]](i[[SZ]] {{[^,]+}}) @@ -158,21 +150,18 @@ int foo(int n) { // CHECK-DAG: [[CPADDR1:%.+]] = bitcast i8** [[PADDR1]] to i[[SZ]]* // CHECK-DAG: store i[[SZ]] [[BP1:%[^,]+]], i[[SZ]]* [[CBPADDR1]] // CHECK-DAG: store i[[SZ]] [[P1:%[^,]+]], i[[SZ]]* [[CPADDR1]] - // CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align 4 - // CHECK-NEXT: br label %[[IFEND:.+]] - - // CHECK: [[IFELSE]] - // CHECK: store i32 -1, i32* [[RHV]], align 4 - // CHECK-NEXT: br label %[[IFEND:.+]] - - // CHECK: [[IFEND]] - // CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align 4 - // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 + // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] // CHECK: [[FAIL]] // CHECK: call void [[HVT3:@.+]]({{[^,]+}}, {{[^,]+}}) // CHECK-NEXT: br label %[[END]] // CHECK: [[END]] + // CHECK-NEXT: br label %[[IFEND:.+]] + // CHECK: [[IFELSE]] + // CHECK: call void [[HVT3]]({{[^,]+}}, {{[^,]+}}) + // CHECK-NEXT: br label %[[IFEND]] + + // CHECK: [[IFEND]] #pragma omp target if(n>10) { a += 1; @@ -285,15 +274,18 @@ int foo(int n) { // CHECK-DAG: store [[TT]]* %{{.+}}, [[TT]]** [[CPADDR8]] // CHECK-DAG: store i[[SZ]] {{12|16}}, i[[SZ]]* [[SADDR8]] - // CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align 4 - // CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align 4 - // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 - // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]] - + // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 + // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] // CHECK: [[FAIL]] // CHECK: call void [[HVT4:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}) // CHECK-NEXT: br label %[[END]] // CHECK: [[END]] + // CHECK-NEXT: br label %[[IFEND:.+]] + // CHECK: [[IFELSE]] + // CHECK: call void [[HVT4]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}) + // CHECK-NEXT: br label %[[IFEND]] + + // CHECK: [[IFEND]] #pragma omp target if(n>20) { a += 1; @@ -521,15 +513,18 @@ int bar(int n){ // CHECK-DAG: store i16* %{{.+}}, i16** [[CPADDR4]] // CHECK-DAG: store i[[SZ]] [[CSIZE]], i[[SZ]]* [[SADDR4]] -// CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align 4 -// CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align 4 -// CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 -// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]] - +// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 +// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] // CHECK: [[FAIL]] // CHECK: call void [[HVT7:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}) // CHECK-NEXT: br label %[[END]] // CHECK: [[END]] +// CHECK-NEXT: br label %[[IFEND:.+]] +// CHECK: [[IFELSE]] +// CHECK: call void [[HVT7]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}) +// CHECK-NEXT: br label %[[IFEND]] + +// CHECK: [[IFEND]] // // CHECK: define {{.*}}[[FSTATIC]] @@ -569,21 +564,18 @@ int bar(int n){ // CHECK-DAG: store [10 x i32]* [[VAL3:%[^,]+]], [10 x i32]** [[CBPADDR3]] // CHECK-DAG: store [10 x i32]* [[VAL3]], [10 x i32]** [[CPADDR3]] -// CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align 4 -// CHECK-NEXT: br label %[[IFEND:.+]] - -// CHECK: [[IFELSE]] -// CHECK: store i32 -1, i32* [[RHV]], align 4 -// CHECK-NEXT: br label %[[IFEND:.+]] - -// CHECK: [[IFEND]] -// CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align 4 -// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 +// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] // CHECK: [[FAIL]] // CHECK: call void [[HVT6:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}) // CHECK-NEXT: br label %[[END]] // CHECK: [[END]] +// CHECK-NEXT: br label %[[IFEND:.+]] +// CHECK: [[IFELSE]] +// CHECK: call void [[HVT6]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}) +// CHECK-NEXT: br label %[[IFEND]] + +// CHECK: [[IFEND]] // // CHECK: define {{.*}}[[FTEMPLATE]] @@ -616,22 +608,18 @@ int bar(int n){ // CHECK-DAG: store [10 x i32]* [[VAL2:%[^,]+]], [10 x i32]** [[CBPADDR2]] // CHECK-DAG: store [10 x i32]* [[VAL2]], [10 x i32]** [[CPADDR2]] -// CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align 4 -// CHECK-NEXT: br label %[[IFEND:.+]] - -// CHECK: [[IFELSE]] -// CHECK: store i32 -1, i32* [[RHV]], align 4 -// CHECK-NEXT: br label %[[IFEND:.+]] - -// CHECK: [[IFEND]] -// CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align 4 -// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 +// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] // CHECK: [[FAIL]] // CHECK: call void [[HVT5:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}) // CHECK-NEXT: br label %[[END]] // CHECK: [[END]] +// CHECK-NEXT: br label %[[IFEND:.+]] +// CHECK: [[IFELSE]] +// CHECK: call void [[HVT5]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}) +// CHECK-NEXT: br label %[[IFEND]] +// CHECK: [[IFEND]] // Check that the offloading functions are emitted and that the arguments are diff --git a/clang/test/OpenMP/target_firstprivate_codegen.cpp b/clang/test/OpenMP/target_firstprivate_codegen.cpp index 64f5225..647b7b6 100644 --- a/clang/test/OpenMP/target_firstprivate_codegen.cpp +++ b/clang/test/OpenMP/target_firstprivate_codegen.cpp @@ -68,7 +68,6 @@ int foo(int n, double *ptr) { // CHECK: [[C:%.+]] = alloca [5 x [10 x double]], // CHECK: [[D:%.+]] = alloca [[TT]], // CHECK: [[ACAST:%.+]] = alloca i{{[0-9]+}}, - // CHECK: {{.+}} = alloca i{{[0-9]+}}, // CHECK: [[BASE_PTR_ARR:%.+]] = alloca [1 x i8*], // CHECK: [[PTR_ARR:%.+]] = alloca [1 x i8*], // CHECK: [[A2CAST:%.+]] = alloca i{{[0-9]+}}, diff --git a/clang/test/OpenMP/target_parallel_codegen.cpp b/clang/test/OpenMP/target_parallel_codegen.cpp index c94924e..17318d5 100644 --- a/clang/test/OpenMP/target_parallel_codegen.cpp +++ b/clang/test/OpenMP/target_parallel_codegen.cpp @@ -94,9 +94,7 @@ int foo(int n) { TT d; // CHECK: [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 0, i8** null, i8** null, i[[SZ]]* null, i32* null, i32 1, i32 0) - // CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align 4 - // CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align 4 - // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 + // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]] // CHECK: [[FAIL]] // CHECK: call void [[HVT0:@.+]]() @@ -106,10 +104,6 @@ int foo(int n) { { } - // CHECK: store i32 0, i32* [[RHV:%.+]], align 4 - // CHECK: store i32 -1, i32* [[RHV]], align 4 - // CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align 4 - // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 // CHECK: call void [[HVT1:@.+]](i[[SZ]] {{[^,]+}}) #pragma omp target parallel if(target: 0) { @@ -126,9 +120,7 @@ int foo(int n) { // CHECK-DAG: store i[[SZ]] [[VAL0:%.+]], i[[SZ]]* [[CBPADDR0]], // CHECK-DAG: store i[[SZ]] [[VAL0]], i[[SZ]]* [[CPADDR0]], - // CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align 4 - // CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align 4 - // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 + // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]] // CHECK: [[FAIL]] // CHECK: call void [[HVT2:@.+]](i[[SZ]] {{[^,]+}}) @@ -159,21 +151,18 @@ int foo(int n) { // CHECK-DAG: [[CPADDR1:%.+]] = bitcast i8** [[PADDR1]] to i[[SZ]]* // CHECK-DAG: store i[[SZ]] [[VAL1:%.+]], i[[SZ]]* [[CBPADDR1]], // CHECK-DAG: store i[[SZ]] [[VAL1]], i[[SZ]]* [[CPADDR1]], - // CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align 4 - // CHECK-NEXT: br label %[[IFEND:.+]] - - // CHECK: [[IFELSE]] - // CHECK: store i32 -1, i32* [[RHV]], align 4 - // CHECK-NEXT: br label %[[IFEND:.+]] - - // CHECK: [[IFEND]] - // CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align 4 - // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 + // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] // CHECK: [[FAIL]] // CHECK: call void [[HVT3:@.+]]({{[^,]+}}, {{[^,]+}}) // CHECK-NEXT: br label %[[END]] // CHECK: [[END]] + // CHECK-NEXT: br label %[[IFEND:.+]] + // CHECK: [[IFELSE]] + // CHECK: call void [[HVT3]]({{[^,]+}}, {{[^,]+}}) + // CHECK-NEXT: br label %[[IFEND]] + // CHECK: [[IFEND]] + #pragma omp target parallel if(target: n>10) { a += 1; @@ -286,9 +275,7 @@ int foo(int n) { // CHECK-DAG: [[CPADDR8]] = bitcast i8** {{%[^,]+}} to [[TT]]** // CHECK-DAG: store i[[SZ]] {{12|16}}, i[[SZ]]* {{%[^,]+}} - // CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align 4 - // CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align 4 - // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 + // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]] // CHECK: [[FAIL]] @@ -582,9 +569,7 @@ int bar(int n){ // CHECK-DAG: [[CPADDR4]] = bitcast i8** {{%[^,]+}} to i16** // CHECK-DAG: store i[[SZ]] [[CSIZE]], i[[SZ]]* {{%[^,]+}} -// CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align 4 -// CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align 4 -// CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 +// CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]] // CHECK: [[FAIL]] @@ -630,21 +615,17 @@ int bar(int n){ // CHECK-DAG: store [10 x i32]* [[VAL3:%.+]], [10 x i32]** [[CBPADDR3]], // CHECK-DAG: store [10 x i32]* [[VAL3]], [10 x i32]** [[CPADDR3]], -// CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align 4 -// CHECK-NEXT: br label %[[IFEND:.+]] - -// CHECK: [[IFELSE]] -// CHECK: store i32 -1, i32* [[RHV]], align 4 -// CHECK-NEXT: br label %[[IFEND:.+]] - -// CHECK: [[IFEND]] -// CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align 4 -// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 +// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] // CHECK: [[FAIL]] // CHECK: call void [[HVT6:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}) // CHECK-NEXT: br label %[[END]] // CHECK: [[END]] +// CHECK-NEXT: br label %[[IFEND:.+]] +// CHECK: [[IFELSE]] +// CHECK: call void [[HVT6]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}) +// CHECK-NEXT: br label %[[IFEND]] +// CHECK: [[IFEND]] // // CHECK: define {{.*}}[[FTEMPLATE]] @@ -677,23 +658,17 @@ int bar(int n){ // CHECK-DAG: store [10 x i32]* [[VAL2:%.+]], [10 x i32]** [[CBPADDR2]], // CHECK-DAG: store [10 x i32]* [[VAL2]], [10 x i32]** [[CPADDR2]], -// CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align 4 -// CHECK-NEXT: br label %[[IFEND:.+]] - -// CHECK: [[IFELSE]] -// CHECK: store i32 -1, i32* [[RHV]], align 4 -// CHECK-NEXT: br label %[[IFEND:.+]] - -// CHECK: [[IFEND]] -// CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align 4 -// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 +// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] // CHECK: [[FAIL]] // CHECK: call void [[HVT5:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}) // CHECK-NEXT: br label %[[END]] // CHECK: [[END]] - - +// CHECK-NEXT: br label %[[IFEND:.+]] +// CHECK: [[IFELSE]] +// CHECK: call void [[HVT:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}) +// CHECK-NEXT: br label %[[IFEND]] +// CHECK: [[IFEND]] // Check that the offloading functions are emitted and that the arguments are // correct and loaded correctly for the target regions of the callees of bar(). diff --git a/clang/test/OpenMP/target_parallel_if_codegen.cpp b/clang/test/OpenMP/target_parallel_if_codegen.cpp index 0aebcc2..215d8f9 100644 --- a/clang/test/OpenMP/target_parallel_if_codegen.cpp +++ b/clang/test/OpenMP/target_parallel_if_codegen.cpp @@ -130,8 +130,6 @@ int bar(int n){ return a; } - - // // CHECK: define {{.*}}[[FS1]]([[S1]]* {{%.+}}, i32 {{[^%]*}}[[PARM:%.+]]) // @@ -148,9 +146,7 @@ int bar(int n){ // CHECK: [[ARG:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR]], align // // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 3, {{.*}}, i32 1, i32 0) -// CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align -// CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align -// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 +// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] // // CHECK: [[FAIL]] @@ -176,27 +172,17 @@ int bar(int n){ // // CHECK: [[IF_THEN]] // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 2, {{.*}}, i32 1, i32 0) -// CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align -// CHECK: br label {{%?}}[[END:.+]] -// -// CHECK: [[IF_ELSE]] -// CHECK: store i32 -1, i32* [[RHV]], align -// CHECK: br label {{%?}}[[END]] -// -// CHECK: [[END]] -// CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align -// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 -// CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] -// +// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 +// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] // CHECK: [[FAIL]] // CHECK: call void [[HVT2:@.+]]([[S1]]* {{%.+}}, i[[SZ]] [[ARG]]) -// CHECK: br label {{%?}}[[END]] +// CHECK-NEXT: br label %[[END]] // CHECK: [[END]] - - - - - +// CHECK-NEXT: br label %[[IFEND:.+]] +// CHECK: [[IF_ELSE]] +// CHECK: call void [[HVT2]]([[S1]]* {{%.+}}, i[[SZ]] [[ARG]]) +// CHECK-NEXT: br label %[[IFEND]] +// CHECK: [[IFEND]] // // CHECK: define {{.*}}[[FSTATIC]](i32 {{[^%]*}}[[PARM:%.+]]) @@ -218,22 +204,17 @@ int bar(int n){ // // CHECK: [[IF_THEN]] // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 1, {{.*}}, i32 1, i32 0) -// CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align -// CHECK: br label {{%?}}[[END:.+]] -// -// CHECK: [[IF_ELSE]] -// CHECK: store i32 -1, i32* [[RHV]], align -// CHECK: br label {{%?}}[[END]] -// -// CHECK: [[END]] -// CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align -// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 -// CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] -// +// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 +// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] // CHECK: [[FAIL]] // CHECK: call void [[HVT3:@.+]](i[[SZ]] [[ARG]]) -// CHECK: br label {{%?}}[[END]] +// CHECK-NEXT: br label %[[END]] // CHECK: [[END]] +// CHECK-NEXT: br label %[[IFEND:.+]] +// CHECK: [[IF_ELSE]] +// CHECK: call void [[HVT3]](i[[SZ]] [[ARG]]) +// CHECK-NEXT: br label %[[IFEND]] +// CHECK: [[IFEND]] // // // @@ -244,22 +225,17 @@ int bar(int n){ // // CHECK: [[IF_THEN]] // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 0, {{.*}}, i32 1, i32 0) -// CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align -// CHECK: br label {{%?}}[[END:.+]] -// -// CHECK: [[IF_ELSE]] -// CHECK: store i32 -1, i32* [[RHV]], align -// CHECK: br label {{%?}}[[END]] -// -// CHECK: [[END]] -// CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align -// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 -// CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] -// +// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 +// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] // CHECK: [[FAIL]] // CHECK: call void [[HVT4:@.+]]() -// CHECK: br label {{%?}}[[END]] +// CHECK-NEXT: br label %[[END]] // CHECK: [[END]] +// CHECK-NEXT: br label %[[IFEND:.+]] +// CHECK: [[IF_ELSE]] +// CHECK: call void [[HVT4]]() +// CHECK-NEXT: br label %[[IFEND]] +// CHECK: [[IFEND]] @@ -270,9 +246,7 @@ int bar(int n){ // CHECK: define {{.*}}[[FTEMPLATE]] // // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 1, {{.*}}, i32 1, i32 0) -// CHECK-NEXT: store i32 [[RET]], i32* [[RHV:%.+]], align -// CHECK-NEXT: [[RET2:%.+]] = load i32, i32* [[RHV]], align -// CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 +// CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] // // CHECK: [[FAIL]] @@ -284,9 +258,7 @@ int bar(int n){ // // // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 2, {{.*}}, i32 1, i32 0) -// CHECK-NEXT: store i32 [[RET]], i32* [[RHV:%.+]], align -// CHECK-NEXT: [[RET2:%.+]] = load i32, i32* [[RHV]], align -// CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 +// CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] // // CHECK: [[FAIL]] diff --git a/clang/test/OpenMP/target_parallel_num_threads_codegen.cpp b/clang/test/OpenMP/target_parallel_num_threads_codegen.cpp index 225c2dd..e2bddae 100644 --- a/clang/test/OpenMP/target_parallel_num_threads_codegen.cpp +++ b/clang/test/OpenMP/target_parallel_num_threads_codegen.cpp @@ -148,9 +148,7 @@ int bar(int n){ // CHECK: [[THREADS:%.+]] = load i32, i32* [[CAPE_ADDR]], align // // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 3, {{.*}}, i32 1, i32 [[THREADS]]) -// CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align -// CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align -// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 +// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] // // CHECK: [[FAIL]] @@ -161,9 +159,7 @@ int bar(int n){ // // // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 1, {{.+}}, i32 1, i32 1024) -// CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align -// CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align -// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 +// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] // // CHECK: [[FAIL]] @@ -191,9 +187,7 @@ int bar(int n){ // CHECK: [[THREADS:%.+]] = load i32, i32* [[CAPE_ADDR]], align // // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 1, {{.*}}, i32 1, i32 [[THREADS]]) -// CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align -// CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align -// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 +// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] // // CHECK: [[FAIL]] @@ -214,9 +208,7 @@ int bar(int n){ // CHECK: [[THREADS:%.+]] = load i32, i32* [[CAPE_ADDR]], align // // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 1, {{.*}}, i32 1, i32 [[THREADS]]) -// CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align -// CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align -// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 +// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] // // CHECK: [[FAIL]] @@ -234,9 +226,7 @@ int bar(int n){ // CHECK: define {{.*}}[[FTEMPLATE]] // // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 0, {{.*}}, i32 1, i32 20) -// CHECK-NEXT: store i32 [[RET]], i32* [[RHV:%.+]], align -// CHECK-NEXT: [[RET2:%.+]] = load i32, i32* [[RHV]], align -// CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 +// CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] // // CHECK: [[FAIL]] @@ -258,9 +248,7 @@ int bar(int n){ // CHECK: [[THREADS:%.+]] = sext i16 [[T]] to i32 // // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 3, {{.*}}, i32 1, i32 [[THREADS]]) -// CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align -// CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align -// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 +// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] // // CHECK: [[FAIL]] diff --git a/clang/test/OpenMP/target_teams_codegen.cpp b/clang/test/OpenMP/target_teams_codegen.cpp index 9def589..978879d 100644 --- a/clang/test/OpenMP/target_teams_codegen.cpp +++ b/clang/test/OpenMP/target_teams_codegen.cpp @@ -96,9 +96,7 @@ int foo(int n) { TT d; // CHECK: [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 0, i8** null, i8** null, i[[SZ]]* null, i32* null, i32 0, i32 0) - // CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align 4 - // CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align 4 - // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 + // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]] // CHECK: [[FAIL]] // CHECK: call void [[HVT0:@.+]]() @@ -108,10 +106,6 @@ int foo(int n) { { } - // CHECK: store i32 0, i32* [[RHV:%.+]], align 4 - // CHECK: store i32 -1, i32* [[RHV]], align 4 - // CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align 4 - // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 // CHECK: call void [[HVT1:@.+]](i[[SZ]] {{[^,]+}}) #pragma omp target teams if(target: 0) { @@ -128,9 +122,7 @@ int foo(int n) { // CHECK-DAG: store i[[SZ]] %{{.+}}, i[[SZ]]* [[CBPADDR0]] // CHECK-DAG: store i[[SZ]] %{{.+}}, i[[SZ]]* [[CPADDR0]] - // CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align 4 - // CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align 4 - // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 + // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]] // CHECK: [[FAIL]] // CHECK: call void [[HVT2:@.+]](i[[SZ]] {{[^,]+}}) @@ -161,21 +153,17 @@ int foo(int n) { // CHECK-DAG: [[CPADDR1:%.+]] = bitcast i8** [[PADDR1]] to i[[SZ]]* // CHECK-DAG: store i[[SZ]] %{{.+}}, i[[SZ]]* [[CBPADDR1]] // CHECK-DAG: store i[[SZ]] %{{.+}}, i[[SZ]]* [[CPADDR1]] - // CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align 4 - // CHECK-NEXT: br label %[[IFEND:.+]] - - // CHECK: [[IFELSE]] - // CHECK: store i32 -1, i32* [[RHV]], align 4 - // CHECK-NEXT: br label %[[IFEND:.+]] - - // CHECK: [[IFEND]] - // CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align 4 - // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 + // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] // CHECK: [[FAIL]] // CHECK: call void [[HVT3:@.+]]({{[^,]+}}, {{[^,]+}}) // CHECK-NEXT: br label %[[END]] // CHECK: [[END]] + // CHECK-NEXT: br label %[[IFEND:.+]] + // CHECK: [[IFELSE]] + // CHECK: call void [[HVT3]]({{[^,]+}}, {{[^,]+}}) + // CHECK-NEXT: br label %[[IFEND]] + // CHECK: [[IFEND]] #pragma omp target teams if(target: n>10) { a += 1; @@ -289,9 +277,7 @@ int foo(int n) { // CHECK-DAG: [[CPADDR8]] = bitcast i8** {{%[^,]+}} to [[TT]]** // CHECK-DAG: store i[[SZ]] {{12|16}}, i[[SZ]]* {{%[^,]+}} - // CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align 4 - // CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align 4 - // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 + // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]] // CHECK: [[FAIL]] @@ -596,9 +582,7 @@ int bar(int n){ // CHECK-DAG: [[CPADDR4]] = bitcast i8** {{%[^,]+}} to i16** // CHECK-DAG: store i[[SZ]] [[CSIZE]], i[[SZ]]* {{%[^,]+}} -// CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align 4 -// CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align 4 -// CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 +// CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]] // CHECK: [[FAIL]] @@ -644,21 +628,17 @@ int bar(int n){ // CHECK-DAG: [[CBPADDR3]] = bitcast i8** {{%[^,]+}} to [10 x i32]** // CHECK-DAG: [[CPADDR3]] = bitcast i8** {{%[^,]+}} to [10 x i32]** -// CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align 4 -// CHECK-NEXT: br label %[[IFEND:.+]] - -// CHECK: [[IFELSE]] -// CHECK: store i32 -1, i32* [[RHV]], align 4 -// CHECK-NEXT: br label %[[IFEND:.+]] - -// CHECK: [[IFEND]] -// CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align 4 -// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 +// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] // CHECK: [[FAIL]] // CHECK: call void [[HVT6:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}) // CHECK-NEXT: br label %[[END]] // CHECK: [[END]] +// CHECK-NEXT: br label %[[IFEND:.+]] +// CHECK: [[IFELSE]] +// CHECK: call void [[HVT6]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}) +// CHECK-NEXT: br label %[[IFEND]] +// CHECK: [[IFEND]] // // CHECK: define {{.*}}[[FTEMPLATE]] @@ -691,21 +671,17 @@ int bar(int n){ // CHECK-DAG: [[CBPADDR2]] = bitcast i8** {{%[^,]+}} to [10 x i32]** // CHECK-DAG: [[CPADDR2]] = bitcast i8** {{%[^,]+}} to [10 x i32]** -// CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align 4 -// CHECK-NEXT: br label %[[IFEND:.+]] - -// CHECK: [[IFELSE]] -// CHECK: store i32 -1, i32* [[RHV]], align 4 -// CHECK-NEXT: br label %[[IFEND:.+]] - -// CHECK: [[IFEND]] -// CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align 4 -// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 +// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] // CHECK: [[FAIL]] // CHECK: call void [[HVT5:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}) // CHECK-NEXT: br label %[[END]] // CHECK: [[END]] +// CHECK-NEXT: br label %[[IFEND:.+]] +// CHECK: [[IFELSE]] +// CHECK: call void [[HVT5]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}) +// CHECK-NEXT: br label %[[IFEND]] +// CHECK: [[IFEND]] diff --git a/clang/test/OpenMP/target_teams_num_teams_codegen.cpp b/clang/test/OpenMP/target_teams_num_teams_codegen.cpp index ad36265..068f76e 100644 --- a/clang/test/OpenMP/target_teams_num_teams_codegen.cpp +++ b/clang/test/OpenMP/target_teams_num_teams_codegen.cpp @@ -148,9 +148,7 @@ int bar(int n){ // CHECK: [[TEAMS:%.+]] = load i32, i32* [[CAPE_ADDR]], align // // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 3, {{.*}}, i32 [[TEAMS]], i32 0) -// CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align -// CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align -// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 +// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] // // CHECK: [[FAIL]] @@ -161,9 +159,7 @@ int bar(int n){ // // // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 1, {{.+}}, i32 1024, i32 0) -// CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align -// CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align -// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 +// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] // // CHECK: [[FAIL]] @@ -191,9 +187,7 @@ int bar(int n){ // CHECK: [[TEAMS:%.+]] = load i32, i32* [[CAPE_ADDR]], align // // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 1, {{.*}}, i32 [[TEAMS]], i32 0) -// CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align -// CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align -// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 +// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] // // CHECK: [[FAIL]] @@ -214,9 +208,7 @@ int bar(int n){ // CHECK: [[TEAMS:%.+]] = load i32, i32* [[CAPE_ADDR]], align // // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 1, {{.*}}, i32 [[TEAMS]], i32 0) -// CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align -// CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align -// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 +// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] // // CHECK: [[FAIL]] @@ -234,9 +226,7 @@ int bar(int n){ // CHECK: define {{.*}}[[FTEMPLATE]] // // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 0, {{.*}}, i32 20, i32 0) -// CHECK-NEXT: store i32 [[RET]], i32* [[RHV:%.+]], align -// CHECK-NEXT: [[RET2:%.+]] = load i32, i32* [[RHV]], align -// CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 +// CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] // // CHECK: [[FAIL]] @@ -258,9 +248,7 @@ int bar(int n){ // CHECK: [[TEAMS:%.+]] = sext i16 [[T]] to i32 // // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 3, {{.*}}, i32 [[TEAMS]], i32 0) -// CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align -// CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align -// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 +// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] // // CHECK: [[FAIL]] diff --git a/clang/test/OpenMP/target_teams_thread_limit_codegen.cpp b/clang/test/OpenMP/target_teams_thread_limit_codegen.cpp index d89f515..c56d142 100644 --- a/clang/test/OpenMP/target_teams_thread_limit_codegen.cpp +++ b/clang/test/OpenMP/target_teams_thread_limit_codegen.cpp @@ -148,9 +148,7 @@ int bar(int n){ // CHECK: [[TL:%.+]] = load i32, i32* [[CAPE_ADDR]], align // // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 3, {{.*}}, i32 0, i32 [[TL]]) -// CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align -// CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align -// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 +// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] // // CHECK: [[FAIL]] @@ -161,9 +159,7 @@ int bar(int n){ // // // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 1, {{.+}}, i32 0, i32 1024) -// CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align -// CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align -// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 +// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] // // CHECK: [[FAIL]] @@ -200,9 +196,7 @@ int bar(int n){ // CHECK: [[TL:%.+]] = load i32, i32* [[CAPE_ADDR2]], align // // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 2, {{.*}}, i32 [[TEAMS]], i32 [[TL]]) -// CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align -// CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align -// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 +// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] // // CHECK: [[FAIL]] @@ -223,9 +217,7 @@ int bar(int n){ // CHECK: [[TL:%.+]] = load i32, i32* [[CAPE_ADDR]], align // // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 1, {{.*}}, i32 0, i32 [[TL]]) -// CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align -// CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align -// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 +// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] // // CHECK: [[FAIL]] @@ -243,9 +235,7 @@ int bar(int n){ // CHECK: define {{.*}}[[FTEMPLATE]] // // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 0, {{.*}}, i32 0, i32 20) -// CHECK-NEXT: store i32 [[RET]], i32* [[RHV:%.+]], align -// CHECK-NEXT: [[RET2:%.+]] = load i32, i32* [[RHV]], align -// CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 +// CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] // // CHECK: [[FAIL]] @@ -267,9 +257,7 @@ int bar(int n){ // CHECK: [[TEAMS:%.+]] = sext i16 [[T]] to i32 // // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 3, {{.*}}, i32 [[TEAMS]], i32 1024) -// CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align -// CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align -// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 +// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] // // CHECK: [[FAIL]]