Author: abataev Date: Thu Dec 7 11:49:28 2017 New Revision: 320078 URL: http://llvm.org/viewvc/llvm-project?rev=320078&view=rev Log: [OPENMP] Do not capture private variables in the target regions.
Private variables are completely redefined in the outlined regions, so we don't need to capture them. Patch adds this behavior to the target-based regions. Modified: cfe/trunk/lib/Sema/SemaExpr.cpp cfe/trunk/test/OpenMP/teams_distribute_parallel_for_private_codegen.cpp cfe/trunk/test/OpenMP/teams_distribute_parallel_for_simd_private_codegen.cpp cfe/trunk/test/OpenMP/teams_distribute_private_codegen.cpp cfe/trunk/test/OpenMP/teams_distribute_simd_private_codegen.cpp cfe/trunk/test/OpenMP/teams_private_codegen.cpp Modified: cfe/trunk/lib/Sema/SemaExpr.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaExpr.cpp?rev=320078&r1=320077&r2=320078&view=diff ============================================================================== --- cfe/trunk/lib/Sema/SemaExpr.cpp (original) +++ cfe/trunk/lib/Sema/SemaExpr.cpp Thu Dec 7 11:49:28 2017 @@ -14619,14 +14619,16 @@ bool Sema::tryCaptureVariable( // just break here. Similarly, global variables that are captured in a // target region should not be captured outside the scope of the region. if (RSI->CapRegionKind == CR_OpenMP) { - auto IsTargetCap = isOpenMPTargetCapturedDecl(Var, RSI->OpenMPLevel); + bool IsOpenMPPrivateDecl = isOpenMPPrivateDecl(Var, RSI->OpenMPLevel); + auto IsTargetCap = !IsOpenMPPrivateDecl && + isOpenMPTargetCapturedDecl(Var, RSI->OpenMPLevel); // When we detect target captures we are looking from inside the // target region, therefore we need to propagate the capture from the // enclosing region. Therefore, the capture is not initially nested. if (IsTargetCap) FunctionScopesIndex--; - if (IsTargetCap || isOpenMPPrivateDecl(Var, RSI->OpenMPLevel)) { + if (IsTargetCap || IsOpenMPPrivateDecl) { Nested = !IsTargetCap; DeclRefType = DeclRefType.getUnqualifiedType(); CaptureType = Context.getLValueReferenceType(DeclRefType); Modified: cfe/trunk/test/OpenMP/teams_distribute_parallel_for_private_codegen.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/teams_distribute_parallel_for_private_codegen.cpp?rev=320078&r1=320077&r2=320078&view=diff ============================================================================== --- cfe/trunk/test/OpenMP/teams_distribute_parallel_for_private_codegen.cpp (original) +++ cfe/trunk/test/OpenMP/teams_distribute_parallel_for_private_codegen.cpp Thu Dec 7 11:49:28 2017 @@ -72,13 +72,13 @@ int main() { // LAMBDA: call void [[OUTER_LAMBDA:@.+]]( [&]() { // LAMBDA: define{{.*}} internal{{.*}} void [[OUTER_LAMBDA]]( - // LAMBDA: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 3, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0) + // LAMBDA: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 1, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0) // LAMBDA: call void @[[LOFFL1:.+]]( // LAMBDA: ret #pragma omp target #pragma omp teams distribute parallel for private(g, g1, sivar) for (int i = 0; i < 2; ++i) { - // LAMBDA: define{{.*}} internal{{.*}} void @[[LOFFL1]](i{{64|32}} {{%.+}}, i{{64|32}} {{%.+}}) + // LAMBDA: define{{.*}} internal{{.*}} void @[[LOFFL1]](i{{64|32}} {{%.+}}) // LAMBDA: call void {{.+}} @__kmpc_fork_teams({{.+}}, i32 0, {{.+}} @[[LOUTL1:.+]] to {{.+}}) // LAMBDA: ret void @@ -164,12 +164,12 @@ int main() { } // CHECK: define {{.*}}i{{[0-9]+}} @main() -// CHECK: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 5, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0) -// CHECK: call void @[[OFFL1:.+]](i{{64|32}} %{{.+}}) +// CHECK: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 0, i8** null, i8** null, i{{64|32}}* null, i64* null, i32 0, i32 0) +// CHECK: call void @[[OFFL1:.+]]() // CHECK: {{%.+}} = call{{.*}} i32 @[[TMAIN_INT:.+]]() // CHECK: ret -// CHECK: define{{.*}} void @[[OFFL1]]({{.+}}) +// CHECK: define{{.*}} void @[[OFFL1]]() // CHECK: call void {{.+}} @__kmpc_fork_teams({{.+}}, i32 0, {{.+}} @[[OUTL1:.+]] to {{.+}}) // CHECK: ret void Modified: cfe/trunk/test/OpenMP/teams_distribute_parallel_for_simd_private_codegen.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/teams_distribute_parallel_for_simd_private_codegen.cpp?rev=320078&r1=320077&r2=320078&view=diff ============================================================================== --- cfe/trunk/test/OpenMP/teams_distribute_parallel_for_simd_private_codegen.cpp (original) +++ cfe/trunk/test/OpenMP/teams_distribute_parallel_for_simd_private_codegen.cpp Thu Dec 7 11:49:28 2017 @@ -72,13 +72,13 @@ int main() { // LAMBDA: call void [[OUTER_LAMBDA:@.+]]( [&]() { // LAMBDA: define{{.*}} internal{{.*}} void [[OUTER_LAMBDA]]( - // LAMBDA: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 3, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0) + // LAMBDA: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 1, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0) // LAMBDA: call void @[[LOFFL1:.+]]( // LAMBDA: ret #pragma omp target #pragma omp teams distribute parallel for simd private(g, g1, sivar) for (int i = 0; i < 2; ++i) { - // LAMBDA: define{{.*}} internal{{.*}} void @[[LOFFL1]](i{{64|32}} {{%.+}}, i{{64|32}} {{%.+}}) + // LAMBDA: define{{.*}} internal{{.*}} void @[[LOFFL1]](i{{64|32}} {{%.+}}) // LAMBDA: call void {{.+}} @__kmpc_fork_teams({{.+}}, i32 0, {{.+}} @[[LOUTL1:.+]] to {{.+}}) // LAMBDA: ret void @@ -167,12 +167,12 @@ int main() { } // CHECK: define {{.*}}i{{[0-9]+}} @main() -// CHECK: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 5, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0) -// CHECK: call void @[[OFFL1:.+]](i{{64|32}} %{{.+}}) +// CHECK: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 0, i8** null, i8** null, i{{64|32}}* null, i64* null, i32 0, i32 0) +// CHECK: call void @[[OFFL1:.+]]() // CHECK: {{%.+}} = call{{.*}} i32 @[[TMAIN_INT:.+]]() // CHECK: ret -// CHECK: define{{.*}} void @[[OFFL1]]({{.+}}) +// CHECK: define{{.*}} void @[[OFFL1]]() // CHECK: call void {{.+}} @__kmpc_fork_teams({{.+}}, i32 0, {{.+}} @[[OUTL1:.+]] to {{.+}}) // CHECK: ret void Modified: cfe/trunk/test/OpenMP/teams_distribute_private_codegen.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/teams_distribute_private_codegen.cpp?rev=320078&r1=320077&r2=320078&view=diff ============================================================================== --- cfe/trunk/test/OpenMP/teams_distribute_private_codegen.cpp (original) +++ cfe/trunk/test/OpenMP/teams_distribute_private_codegen.cpp Thu Dec 7 11:49:28 2017 @@ -72,13 +72,13 @@ int main() { // LAMBDA: call void [[OUTER_LAMBDA:@.+]]( [&]() { // LAMBDA: define{{.*}} internal{{.*}} void [[OUTER_LAMBDA]]( - // LAMBDA: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 3, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0) + // LAMBDA: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 1, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0) // LAMBDA: call void @[[LOFFL1:.+]]( // LAMBDA: ret #pragma omp target #pragma omp teams distribute private(g, g1, sivar) for (int i = 0; i < 2; ++i) { - // LAMBDA: define{{.*}} internal{{.*}} void @[[LOFFL1]](i{{64|32}} {{%.+}}, i{{64|32}} {{%.+}}) + // LAMBDA: define{{.*}} internal{{.*}} void @[[LOFFL1]](i{{64|32}} {{%.+}}) // LAMBDA: call void {{.+}} @__kmpc_fork_teams({{.+}}, i32 0, {{.+}} @[[LOUTL1:.+]] to {{.+}}) // LAMBDA: ret void @@ -97,7 +97,6 @@ int main() { // LAMBDA: [[TMP:%.+]] = alloca i{{[0-9]+}}*, // LAMBDA: [[SIVAR_PRIV:%.+]] = alloca i{{[0-9]+}}, // LAMBDA: store{{.+}} [[G1_PRIV]], {{.+}} [[TMP]], - g = 1; g1 = 1; sivar = 2; @@ -142,12 +141,12 @@ int main() { } // CHECK: define {{.*}}i{{[0-9]+}} @main() -// CHECK: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 5, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0) -// CHECK: call void @[[OFFL1:.+]](i{{64|32}} %{{.+}}) +// CHECK: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 0, i8** null, i8** null, i{{64|32}}* null, i64* null, i32 0, i32 0) +// CHECK: call void @[[OFFL1:.+]]() // CHECK: {{%.+}} = call{{.*}} i32 @[[TMAIN_INT:.+]]() // CHECK: ret -// CHECK: define{{.*}} void @[[OFFL1]]({{.+}}) +// CHECK: define{{.*}} void @[[OFFL1]]() // CHECK: call void {{.+}} @__kmpc_fork_teams({{.+}}, i32 0, {{.+}} @[[OUTL1:.+]] to {{.+}}) // CHECK: ret void Modified: cfe/trunk/test/OpenMP/teams_distribute_simd_private_codegen.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/teams_distribute_simd_private_codegen.cpp?rev=320078&r1=320077&r2=320078&view=diff ============================================================================== --- cfe/trunk/test/OpenMP/teams_distribute_simd_private_codegen.cpp (original) +++ cfe/trunk/test/OpenMP/teams_distribute_simd_private_codegen.cpp Thu Dec 7 11:49:28 2017 @@ -72,13 +72,13 @@ int main() { // LAMBDA: call void [[OUTER_LAMBDA:@.+]]( [&]() { // LAMBDA: define{{.*}} internal{{.*}} void [[OUTER_LAMBDA]]( - // LAMBDA: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 3, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0) + // LAMBDA: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 1, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0) // LAMBDA: call void @[[LOFFL1:.+]]( // LAMBDA: ret #pragma omp target #pragma omp teams distribute simd private(g, g1, sivar) for (int i = 0; i < 2; ++i) { - // LAMBDA: define{{.*}} internal{{.*}} void @[[LOFFL1]](i{{64|32}} {{%.+}}, i{{64|32}} {{%.+}}) + // LAMBDA: define{{.*}} internal{{.*}} void @[[LOFFL1]](i{{64|32}} {{%.+}}) // LAMBDA: call void {{.+}} @__kmpc_fork_teams({{.+}}, i32 0, {{.+}} @[[LOUTL1:.+]] to {{.+}}) // LAMBDA: ret void @@ -142,12 +142,12 @@ int main() { } // CHECK: define {{.*}}i{{[0-9]+}} @main() -// CHECK: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 5, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0) -// CHECK: call void @[[OFFL1:.+]](i{{64|32}} %{{.+}}) +// CHECK: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 0, i8** null, i8** null, i{{64|32}}* null, i64* null, i32 0, i32 0) +// CHECK: call void @[[OFFL1:.+]]() // CHECK: {{%.+}} = call{{.*}} i32 @[[TMAIN_INT:.+]]() // CHECK: ret -// CHECK: define{{.*}} void @[[OFFL1]]({{.+}}) +// CHECK: define{{.*}} void @[[OFFL1]]() // CHECK: call void {{.+}} @__kmpc_fork_teams({{.+}}, i32 0, {{.+}} @[[OUTL1:.+]] to {{.+}}) // CHECK: ret void Modified: cfe/trunk/test/OpenMP/teams_private_codegen.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/teams_private_codegen.cpp?rev=320078&r1=320077&r2=320078&view=diff ============================================================================== --- cfe/trunk/test/OpenMP/teams_private_codegen.cpp (original) +++ cfe/trunk/test/OpenMP/teams_private_codegen.cpp Thu Dec 7 11:49:28 2017 @@ -98,7 +98,7 @@ int main() { // lambda and target region in main // LAMBDA: define {{.+}} [[OUTER_LAMBDA]]([[CAP_TY]]* {{.+}}) - // LAMBDA: call void @[[OMP_OFFLOADING:.+]](i{{[0-9]+}}* {{.+}}, i{{[0-9]+}} {{.+}} + // LAMBDA: call void @[[OMP_OFFLOADING:.+]]() // target region in struct constructor // LAMBDA: define{{.*}} void [[ST_CONSTR:@.+]]([[SS_TY]]* %this, @@ -154,7 +154,7 @@ int main() { #pragma omp target #pragma omp teams private(g, sivar) { - // LAMBDA: define{{.+}} @[[OMP_OFFLOADING]](i{{[0-9]+}}* {{.+}} [[G_IN:%.+]], i{{[0-9]+}} [[SIVAR_IN:%.+]] + // LAMBDA: define{{.+}} @[[OMP_OFFLOADING]]() // LAMBDA: call void (%{{.+}}*, i{{[0-9]+}}, void (i{{[0-9]+}}*, i{{[0-9]+}}*, ...)*, ...) @__kmpc_fork_teams(%{{.+}}* @{{.+}}, i{{[0-9]+}} 0, void (i{{[0-9]+}}*, i{{[0-9]+}}*, ...)* bitcast (void (i{{[0-9]+}}*, i{{[0-9]+}}*)* [[OMP_OUTLINED_1:@.+]] to void // LAMBDA: define {{.+}} [[OMP_OUTLINED_1]](i{{[0-9]+}}* {{.+}}, i{{[0-9]+}}* {{.+}} @@ -196,14 +196,13 @@ int main() { // CHECK: define{{.*}} i{{[0-9]+}} @main() // CHECK: [[TEST:%.+]] = alloca [[S_FLOAT_TY]], // CHECK: call {{.*}} [[S_FLOAT_TY_DEF_CONSTR:@.+]]([[S_FLOAT_TY]]* [[TEST]]) -// CHECK: [[OFF_IN:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* {{%.+}}, -// CHECK: call void @[[OMP_OFFLOADING:.+]](i{{[0-9]+}} [[OFF_IN]] +// CHECK: call void @[[OMP_OFFLOADING:.+]]() // CHECK: = call{{.*}} i{{.+}} [[TMAIN_INT:@.+]]() // CHECK: call void [[S_FLOAT_TY_DESTR:@.+]]([[S_FLOAT_TY]]* // CHECK: ret // target region in main function -// CHECK: define{{.+}} @[[OMP_OFFLOADING]](i{{[0-9]+}} +// CHECK: define{{.+}} @[[OMP_OFFLOADING]]() // CHECK: call void (%{{.+}}*, i{{[0-9]+}}, void (i{{[0-9]+}}*, i{{[0-9]+}}*, ...)*, ...) @__kmpc_fork_teams(%{{.+}}* @{{.+}}, i{{[0-9]+}} 0, void (i{{[0-9]+}}*, i{{[0-9]+}}*, ...)* bitcast (void (i{{[0-9]+}}*, i{{[0-9]+}}*)* [[OMP_OUTLINED:@.+]] to void // CHECK: define internal void [[OMP_OUTLINED]](i{{[0-9]+}}* noalias [[GTID_ADDR:%.+]], i{{[0-9]+}}* noalias %{{.+}}) _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits