r320078 - [OPENMP] Do not capture private variables in the target regions.
Alexey Bataev via cfe-commits
cfe-commits at lists.llvm.org
Thu Dec 7 11:49:28 PST 2017
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 %{{.+}})
More information about the cfe-commits
mailing list