[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.

llvm-svn: 347096

2 files changed, 27 insertions, 11 deletions

diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index 59ac0348cd4..c5ac084e17f 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/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<const VarDecl *, FieldDecl *> 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/clang/test/OpenMP/nvptx_lambda_capturing.cpp b/clang/test/OpenMP/nvptx_lambda_capturing.cpp
index 481d4ad7450..f57fef638f4 100644
--- a/clang/test/OpenMP/nvptx_lambda_capturing.cpp
+++ b/clang/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 <typename T>
+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