From a495c6403b5348bc0243393ab5ee467983b3c8eb Mon Sep 17 00:00:00 2001
From: Alexey Bataev <a.bataev@hotmail.com>
Date: Mon, 11 Mar 2019 19:51:42 +0000
Subject: [PATCH] [OPENMP]Fix codegen for declare target link in target
 regions.

If the declare target link global is used in the target region
indirectly (used in the inner parallel, teams, etc. regions), we may
miss this variable and it leads to incorrect codegen.

llvm-svn: 355858
---
 clang/lib/Sema/SemaOpenMP.cpp                     | 48 ++++++++++++++++++++++-
 clang/test/OpenMP/declare_target_link_codegen.cpp | 10 +++--
 2 files changed, 53 insertions(+), 5 deletions(-)

diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index 6c4c4a8..32e493e 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -147,6 +147,9 @@ private:
     /// Reference to the taskgroup task_reduction reference expression.
     Expr *TaskgroupReductionRef = nullptr;
     llvm::DenseSet<QualType> MappedClassesQualTypes;
+    /// List of globals marked as declare target link in this target region
+    /// (isOpenMPTargetExecutionDirective(Directive) == true).
+    llvm::SmallVector<DeclRefExpr *, 4> DeclareTargetLinkVarDecls;
     SharingMapTy(OpenMPDirectiveKind DKind, DeclarationNameInfo Name,
                  Scope *CurScope, SourceLocation Loc)
         : Directive(DKind), DirectiveName(Name), CurScope(CurScope),
@@ -674,6 +677,31 @@ public:
     return StackElem.MappedClassesQualTypes.count(QT) != 0;
   }
 
+  /// Adds global declare target to the parent target region.
+  void addToParentTargetRegionLinkGlobals(DeclRefExpr *E) {
+    assert(*OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(
+               E->getDecl()) == OMPDeclareTargetDeclAttr::MT_Link &&
+           "Expected declare target link global.");
+    if (isStackEmpty())
+      return;
+    auto It = Stack.back().first.rbegin();
+    while (It != Stack.back().first.rend() &&
+           !isOpenMPTargetExecutionDirective(It->Directive))
+      ++It;
+    if (It != Stack.back().first.rend()) {
+      assert(isOpenMPTargetExecutionDirective(It->Directive) &&
+             "Expected target executable directive.");
+      It->DeclareTargetLinkVarDecls.push_back(E);
+    }
+  }
+
+  /// Returns the list of globals with declare target link if current directive
+  /// is target.
+  ArrayRef<DeclRefExpr *> getLinkGlobals() const {
+    assert(isOpenMPTargetExecutionDirective(getCurrentDirective()) &&
+           "Expected target executable directive.");
+    return Stack.back().first.back().DeclareTargetLinkVarDecls;
+  }
 };
 
 bool isImplicitTaskingRegion(OpenMPDirectiveKind DKind) {
@@ -2414,8 +2442,18 @@ public:
       // Define implicit data-sharing attributes for task.
       DVar = Stack->getImplicitDSA(VD, /*FromParent=*/false);
       if (isOpenMPTaskingDirective(DKind) && DVar.CKind != OMPC_shared &&
-          !Stack->isLoopControlVariable(VD).first)
+          !Stack->isLoopControlVariable(VD).first) {
         ImplicitFirstprivate.push_back(E);
+        return;
+      }
+
+      // Store implicitly used globals with declare target link for parent
+      // target.
+      if (!isOpenMPTargetExecutionDirective(DKind) && Res &&
+          *Res == OMPDeclareTargetDeclAttr::MT_Link) {
+        Stack->addToParentTargetRegionLinkGlobals(E);
+        return;
+      }
     }
   }
   void VisitMemberExpr(MemberExpr *E) {
@@ -2573,7 +2611,13 @@ public:
   }
 
   DSAAttrChecker(DSAStackTy *S, Sema &SemaRef, CapturedStmt *CS)
-      : Stack(S), SemaRef(SemaRef), ErrorFound(false), CS(CS) {}
+      : Stack(S), SemaRef(SemaRef), ErrorFound(false), CS(CS) {
+    // Process declare target link variables for the target directives.
+    if (isOpenMPTargetExecutionDirective(S->getCurrentDirective())) {
+      for (DeclRefExpr *E : Stack->getLinkGlobals())
+        Visit(E);
+    }
+  }
 };
 } // namespace
 
diff --git a/clang/test/OpenMP/declare_target_link_codegen.cpp b/clang/test/OpenMP/declare_target_link_codegen.cpp
index 564581c..36c295b 100644
--- a/clang/test/OpenMP/declare_target_link_codegen.cpp
+++ b/clang/test/OpenMP/declare_target_link_codegen.cpp
@@ -37,10 +37,13 @@ int maini1() {
   {
     a = c;
   }
+#pragma omp target
+#pragma omp teams
+  c = a;
   return 0;
 }
 
-// DEVICE: define weak void @__omp_offloading_{{.*}}_{{.*}}maini1{{.*}}_l[[@LINE-7]](i32* dereferenceable{{[^,]*}}
+// DEVICE: define weak void @__omp_offloading_{{.*}}_{{.*}}maini1{{.*}}_l[[@LINE-10]](i32* dereferenceable{{[^,]*}}
 // DEVICE: [[C_REF:%.+]] = load i32*, i32** @c_decl_tgt_link_ptr,
 // DEVICE: [[C:%.+]] = load i32, i32* [[C_REF]],
 // DEVICE: store i32 [[C]], i32* %
@@ -59,9 +62,10 @@ int maini1() {
 // HOST: [[BP0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BASEPTRS]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
 // HOST: [[P0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTRS]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
 // HOST: call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 2, i8** [[BP0]], i8** [[P0]], i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[SIZES]], i{{[0-9]+}} 0, i{{[0-9]+}} 0), i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[MAPTYPES]], i{{[0-9]+}} 0, i{{[0-9]+}} 0))
-// HOST: call void @__omp_offloading_{{.*}}_{{.*}}maini1{{.*}}_l[[@LINE-26]](i32* %{{[^,]+}})
+// HOST: call void @__omp_offloading_{{.*}}_{{.*}}maini1{{.*}}_l[[@LINE-29]](i32* %{{[^,]+}})
+// HOST: call i32 @__tgt_target_teams(i64 -1, i8* @.__omp_offloading_{{.+}}_l40.region_id, i32 2, {{.+}})
 
-// HOST: define internal void @__omp_offloading_{{.*}}_{{.*}}maini1{{.*}}_l[[@LINE-28]](i32* dereferenceable{{.*}})
+// HOST: define internal void @__omp_offloading_{{.*}}_{{.*}}maini1{{.*}}_l[[@LINE-32]](i32* dereferenceable{{.*}})
 // HOST: [[C:%.*]] = load i32, i32* @c,
 // HOST: store i32 [[C]], i32* %
 
-- 
2.7.4