From d0e625611a370ed08b1179e46a45da34a651d0af Mon Sep 17 00:00:00 2001 From: Alexey Bataev Date: Fri, 16 Nov 2018 21:13:33 +0000 Subject: [PATCH] [OPENMP]Fix PR39694: do not capture `this` in non-`this` region. If lambda is used inside of the OpenMP region and captures `this`, we should recapture it in the OpenMP region also. But we should do this only if the OpenMP region is used in the context of the same class, just like the lambda. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@347096 91177308-0d34-0410-b5e6-96231b3b80d8 --- lib/Sema/SemaOpenMP.cpp | 11 +++++++++-- test/OpenMP/nvptx_lambda_capturing.cpp | 27 +++++++++++++++++--------- 2 files changed, 27 insertions(+), 11 deletions(-) diff --git a/lib/Sema/SemaOpenMP.cpp b/lib/Sema/SemaOpenMP.cpp index 59ac0348cd..c5ac084e17 100644 --- a/lib/Sema/SemaOpenMP.cpp +++ b/lib/Sema/SemaOpenMP.cpp @@ -1533,7 +1533,10 @@ VarDecl *Sema::isOpenMPCapturedDecl(ValueDecl *D) { bool SavedForceCaptureByReferenceInTargetExecutable = DSAStack->isForceCaptureByReferenceInTargetExecutable(); DSAStack->setForceCaptureByReferenceInTargetExecutable(/*V=*/true); - if (RD->isLambda()) + if (RD->isLambda()) { + llvm::DenseMap Captures; + FieldDecl *ThisCapture; + RD->getCaptureFields(Captures, ThisCapture); for (const LambdaCapture &LC : RD->captures()) { if (LC.getCaptureKind() == LCK_ByRef) { VarDecl *VD = LC.getCapturedVar(); @@ -1552,9 +1555,13 @@ VarDecl *Sema::isOpenMPCapturedDecl(ValueDecl *D) { OpenMPClauseKind) { return true; })) MarkVariableReferenced(LC.getLocation(), LC.getCapturedVar()); } else if (LC.getCaptureKind() == LCK_This) { - CheckCXXThisCapture(LC.getLocation()); + QualType ThisTy = getCurrentThisType(); + if (!ThisTy.isNull() && + Context.typesAreCompatible(ThisTy, ThisCapture->getType())) + CheckCXXThisCapture(LC.getLocation()); } } + } DSAStack->setForceCaptureByReferenceInTargetExecutable( SavedForceCaptureByReferenceInTargetExecutable); } diff --git a/test/OpenMP/nvptx_lambda_capturing.cpp b/test/OpenMP/nvptx_lambda_capturing.cpp index 481d4ad745..f57fef638f 100644 --- a/test/OpenMP/nvptx_lambda_capturing.cpp +++ b/test/OpenMP/nvptx_lambda_capturing.cpp @@ -20,15 +20,17 @@ // HOST-DAG: = private unnamed_addr constant [3 x i64] [i64 547, i64 673, i64 562949953422096] // HOST-DAG: = private unnamed_addr constant [3 x i64] [i64 4, i64 8, i64 8] // HOST-DAG: = private unnamed_addr constant [3 x i64] [i64 547, i64 673, i64 562949953422096] +// HOST-DAG: = private unnamed_addr constant [2 x i64] [i64 8, i64 8] +// HOST-DAG: = private unnamed_addr constant [2 x i64] [i64 673, i64 281474976711440] // CHECK-DAG: [[S:%.+]] = type { i32 } // CHECK-DAG: [[CAP1:%.+]] = type { [[S]]* } // CHECK-DAG: [[CAP2:%.+]] = type { i32*, i32*, i32*, i32**, i32* } -// CLASS: define internal void @__omp_offloading_{{.*}}_{{.*}}foo{{.*}}_l63_worker() -// CLASS: define weak void @__omp_offloading_{{.*}}_{{.*}}foo{{.*}}_l63([[S]]* {{%.+}}, [[CAP1]]* dereferenceable(8) {{%.+}}) +// CLASS: define internal void @__omp_offloading_{{.*}}_{{.*}}foo{{.*}}_l72_worker() +// CLASS: define weak void @__omp_offloading_{{.*}}_{{.*}}foo{{.*}}_l72([[S]]* {{%.+}}, [[CAP1]]* dereferenceable(8) {{%.+}}) // CLASS-NOT: getelementptr // CLASS: br i1 % -// CLASS: call void @__omp_offloading_{{.*}}_{{.*}}foo{{.*}}_l63_worker() +// CLASS: call void @__omp_offloading_{{.*}}_{{.*}}foo{{.*}}_l72_worker() // CLASS: br label % // CLASS: br i1 % // CLASS: call void @__kmpc_kernel_init( @@ -41,7 +43,7 @@ // CLASS: call i32 [[LAMBDA1:@.+foo.+]]([[CAP1]]* [[L]]) // CLASS: ret void -// CLASS: define weak void @__omp_offloading_{{.+}}foo{{.+}}_l65([[S]]* %{{.+}}, [[CAP1]]* dereferenceable(8) %{{.+}}) +// CLASS: define weak void @__omp_offloading_{{.+}}foo{{.+}}_l74([[S]]* %{{.+}}, [[CAP1]]* dereferenceable(8) %{{.+}}) // CLASS-NOT: getelementptr // CLASS: call void [[PARALLEL:@.+]](i32* %{{.+}}, i32* %{{.+}}, [[S]]* %{{.+}}, [[CAP1]]* %{{.+}}) // CLASS: ret void @@ -56,6 +58,13 @@ // CLASS: call i32 [[LAMBDA1]]([[CAP1]]* [[L]]) // CLASS: ret void +template +int foo(const T &t) { + #pragma omp target parallel + t(); + return 0; +} + struct S { int a = 15; int foo() { @@ -64,15 +73,15 @@ struct S { L(); #pragma omp target parallel L(); - return a; + return a + ::foo(L); } } s; -// FUN: define internal void @__omp_offloading_{{.+}}_main_l125_worker() -// FUN: define weak void @__omp_offloading_{{.+}}_main_l125(i32* dereferenceable(4) %{{.+}}, i32* dereferenceable(4) %{{.+}}, i32* %{{.+}}, i32* dereferenceable(4) %{{.+}}, [[CAP2]]* dereferenceable(40) %{{.+}}, i64 %{{.+}}) +// FUN: define internal void @__omp_offloading_{{.+}}_main_l134_worker() +// FUN: define weak void @__omp_offloading_{{.+}}_main_l134(i32* dereferenceable(4) %{{.+}}, i32* dereferenceable(4) %{{.+}}, i32* %{{.+}}, i32* dereferenceable(4) %{{.+}}, [[CAP2]]* dereferenceable(40) %{{.+}}, i64 %{{.+}}) // FUN-NOT: getelementptr // FUN: br i1 % -// FUN: call void @__omp_offloading_{{.*}}_{{.*}}main{{.*}}_l125_worker() +// FUN: call void @__omp_offloading_{{.*}}_{{.*}}main{{.*}}_l134_worker() // FUN: br label % // FUN: br i1 % // FUN: call void @__kmpc_kernel_init( @@ -93,7 +102,7 @@ struct S { // FUN: call i64 [[LAMBDA2:@.+main.+]]([[CAP2]]* [[L]]) // FUN: ret void -// FUN: define weak void @__omp_offloading_{{.+}}_main_l127(i32* dereferenceable(4) %{{.+}}, i32* dereferenceable(4) %{{.+}} i32* dereferenceable(4) %{{.+}}, i32* %{{.+}}, i32* dereferenceable(4) %{{.+}}, [[CAP2]]* dereferenceable(40) %{{.+}}) +// FUN: define weak void @__omp_offloading_{{.+}}_main_l136(i32* dereferenceable(4) %{{.+}}, i32* dereferenceable(4) %{{.+}} i32* dereferenceable(4) %{{.+}}, i32* %{{.+}}, i32* dereferenceable(4) %{{.+}}, [[CAP2]]* dereferenceable(40) %{{.+}}) // FUN-NOT: getelementptr // FUN: call void [[PARALLEL:@.+]](i32* %{{.+}}, i32* %{{.+}}, i32* %{{.+}}, i32* %{{.+}}, i32* %{{.+}}, i32* %{{.+}}, i32* %{{.+}}, [[CAP2]]* %{{.+}}) // FUN: ret void -- 2.40.0