[flang-commits] [flang] [OpenMP][libomptarget] Add map checks when running under unified shared memory (PR #69005)

Gheorghe-Teodor Bercea via flang-commits flang-commits at lists.llvm.org
Mon Oct 16 08:51:15 PDT 2023


https://github.com/doru1004 updated https://github.com/llvm/llvm-project/pull/69005

>From cb4121c466a0fc357d6ca129bfdd4e7c5e2d11ee Mon Sep 17 00:00:00 2001
From: Doru Bercea <Doru.Bercea at amd.com>
Date: Wed, 16 Nov 2022 17:23:48 -0600
Subject: [PATCH 1/2] Fix declare target implementation to support enter.

---
 clang/include/clang/Basic/Attr.td             |  4 +-
 .../clang/Basic/DiagnosticParseKinds.td       | 12 ++++-
 clang/lib/AST/AttrImpl.cpp                    |  2 +-
 clang/lib/CodeGen/CGExpr.cpp                  | 12 +++--
 clang/lib/CodeGen/CGOpenMPRuntime.cpp         | 24 ++++++---
 clang/lib/CodeGen/CodeGenModule.cpp           |  6 ++-
 clang/lib/Parse/ParseOpenMP.cpp               | 39 ++++++++++----
 clang/lib/Sema/SemaOpenMP.cpp                 | 10 ++--
 .../test/OpenMP/declare_target_ast_print.cpp  | 53 +++++++++++++++++++
 9 files changed, 130 insertions(+), 32 deletions(-)

diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index 16cf932c3760bd3..eaf4a6db3600e07 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -3749,8 +3749,8 @@ def OMPDeclareTargetDecl : InheritableAttr {
   let Documentation = [OMPDeclareTargetDocs];
   let Args = [
     EnumArgument<"MapType", "MapTypeTy",
-                 [ "to", "link" ],
-                 [ "MT_To", "MT_Link" ]>,
+                 [ "to", "enter", "link" ],
+                 [ "MT_To", "MT_Enter", "MT_Link" ]>,
     EnumArgument<"DevType", "DevTypeTy",
                  [ "host", "nohost", "any" ],
                  [ "DT_Host", "DT_NoHost", "DT_Any" ]>,
diff --git a/clang/include/clang/Basic/DiagnosticParseKinds.td b/clang/include/clang/Basic/DiagnosticParseKinds.td
index 674d6bd34fc544f..27cd3da1f191c3d 100644
--- a/clang/include/clang/Basic/DiagnosticParseKinds.td
+++ b/clang/include/clang/Basic/DiagnosticParseKinds.td
@@ -1383,12 +1383,22 @@ def note_omp_assumption_clause_continue_here
     : Note<"the ignored tokens spans until here">;
 def err_omp_declare_target_unexpected_clause: Error<
   "unexpected '%0' clause, only %select{'device_type'|'to' or 'link'|'to', 'link' or 'device_type'|'device_type', 'indirect'|'to', 'link', 'device_type' or 'indirect'}1 clauses expected">;
+def err_omp_declare_target_unexpected_clause_52: Error<
+  "unexpected '%0' clause, only %select{'device_type'|'enter' or 'link'|'enter', 'link' or 'device_type'|'device_type', 'indirect'|'enter', 'link', 'device_type' or 'indirect'}1 clauses expected">;
 def err_omp_begin_declare_target_unexpected_implicit_to_clause: Error<
   "unexpected '(', only 'to', 'link' or 'device_type' clauses expected for 'begin declare target' directive">;
-def err_omp_declare_target_unexpected_clause_after_implicit_to: Error<
+def err_omp_declare_target_wrong_clause_after_implicit_to: Error<
   "unexpected clause after an implicit 'to' clause">;
+def err_omp_declare_target_wrong_clause_after_implicit_enter: Error<
+  "unexpected clause after an implicit 'enter' clause">;
 def err_omp_declare_target_missing_to_or_link_clause: Error<
   "expected at least one %select{'to' or 'link'|'to', 'link' or 'indirect'}0 clause">;
+def err_omp_declare_target_missing_enter_or_link_clause: Error<
+  "expected at least one %select{'enter' or 'link'|'enter', 'link' or 'indirect'}0 clause">;
+def err_omp_declare_target_unexpected_to_clause: Error<
+  "unexpected 'to' clause, use 'enter' instead">;
+def err_omp_declare_target_unexpected_enter_clause: Error<
+  "unexpected 'enter' clause, use 'to' instead">;
 def err_omp_declare_target_multiple : Error<
   "%0 appears multiple times in clauses on the same declare target directive">;
 def err_omp_declare_target_indirect_device_type: Error<
diff --git a/clang/lib/AST/AttrImpl.cpp b/clang/lib/AST/AttrImpl.cpp
index cecbd703ac61e8c..da842f6b190e74d 100644
--- a/clang/lib/AST/AttrImpl.cpp
+++ b/clang/lib/AST/AttrImpl.cpp
@@ -137,7 +137,7 @@ void OMPDeclareTargetDeclAttr::printPrettyPragma(
   // Use fake syntax because it is for testing and debugging purpose only.
   if (getDevType() != DT_Any)
     OS << " device_type(" << ConvertDevTypeTyToStr(getDevType()) << ")";
-  if (getMapType() != MT_To)
+  if (getMapType() != MT_To && getMapType() != MT_Enter)
     OS << ' ' << ConvertMapTypeTyToStr(getMapType());
   if (Expr *E = getIndirectExpr()) {
     OS << " indirect(";
diff --git a/clang/lib/CodeGen/CGExpr.cpp b/clang/lib/CodeGen/CGExpr.cpp
index ee09a8566c3719e..77085ff34fca233 100644
--- a/clang/lib/CodeGen/CGExpr.cpp
+++ b/clang/lib/CodeGen/CGExpr.cpp
@@ -2495,14 +2495,16 @@ static Address emitDeclTargetVarDeclLValue(CodeGenFunction &CGF,
                                            const VarDecl *VD, QualType T) {
   llvm::Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
       OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD);
-  // Return an invalid address if variable is MT_To and unified
-  // memory is not enabled. For all other cases: MT_Link and
-  // MT_To with unified memory, return a valid address.
-  if (!Res || (*Res == OMPDeclareTargetDeclAttr::MT_To &&
+  // Return an invalid address if variable is MT_To (or MT_Enter starting with
+  // OpenMP 5.2) and unified memory is not enabled. For all other cases: MT_Link
+  // and MT_To (or MT_Enter) with unified memory, return a valid address.
+  if (!Res || ((*Res == OMPDeclareTargetDeclAttr::MT_To ||
+                *Res == OMPDeclareTargetDeclAttr::MT_Enter) &&
                !CGF.CGM.getOpenMPRuntime().hasRequiresUnifiedSharedMemory()))
     return Address::invalid();
   assert(((*Res == OMPDeclareTargetDeclAttr::MT_Link) ||
-          (*Res == OMPDeclareTargetDeclAttr::MT_To &&
+          ((*Res == OMPDeclareTargetDeclAttr::MT_To ||
+            *Res == OMPDeclareTargetDeclAttr::MT_Enter) &&
            CGF.CGM.getOpenMPRuntime().hasRequiresUnifiedSharedMemory())) &&
          "Expected link clause OR to clause with unified memory enabled.");
   QualType PtrTy = CGF.getContext().getPointerType(VD->getType());
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 4d226bbacb98f27..35eda4dca3210c1 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -1627,7 +1627,8 @@ Address CGOpenMPRuntime::getAddrOfDeclareTargetVar(const VarDecl *VD) {
   llvm::Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
       OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD);
   if (Res && (*Res == OMPDeclareTargetDeclAttr::MT_Link ||
-              (*Res == OMPDeclareTargetDeclAttr::MT_To &&
+              ((*Res == OMPDeclareTargetDeclAttr::MT_To ||
+                *Res == OMPDeclareTargetDeclAttr::MT_Enter) &&
                HasRequiresUnifiedSharedMemory))) {
     SmallString<64> PtrName;
     {
@@ -1840,7 +1841,8 @@ bool CGOpenMPRuntime::emitDeclareTargetVarDefinition(const VarDecl *VD,
   Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
       OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD);
   if (!Res || *Res == OMPDeclareTargetDeclAttr::MT_Link ||
-      (*Res == OMPDeclareTargetDeclAttr::MT_To &&
+      ((*Res == OMPDeclareTargetDeclAttr::MT_To ||
+        *Res == OMPDeclareTargetDeclAttr::MT_Enter) &&
        HasRequiresUnifiedSharedMemory))
     return CGM.getLangOpts().OpenMPIsDevice;
   VD = VD->getDefinition(CGM.getContext());
@@ -7461,7 +7463,8 @@ class MappableExprsHandler {
         if (llvm::Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
                 OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD)) {
           if ((*Res == OMPDeclareTargetDeclAttr::MT_Link) ||
-              (*Res == OMPDeclareTargetDeclAttr::MT_To &&
+              ((*Res == OMPDeclareTargetDeclAttr::MT_To ||
+                *Res == OMPDeclareTargetDeclAttr::MT_Enter) &&
                CGF.CGM.getOpenMPRuntime().hasRequiresUnifiedSharedMemory())) {
             RequiresReference = true;
             BP = CGF.CGM.getOpenMPRuntime().getAddrOfDeclareTargetVar(VD);
@@ -10325,7 +10328,8 @@ bool CGOpenMPRuntime::emitTargetGlobalVariable(GlobalDecl GD) {
       OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(
           cast<VarDecl>(GD.getDecl()));
   if (!Res || *Res == OMPDeclareTargetDeclAttr::MT_Link ||
-      (*Res == OMPDeclareTargetDeclAttr::MT_To &&
+      ((*Res == OMPDeclareTargetDeclAttr::MT_To ||
+        *Res == OMPDeclareTargetDeclAttr::MT_Enter) &&
        HasRequiresUnifiedSharedMemory)) {
     DeferredGlobalVariables.insert(cast<VarDecl>(GD.getDecl()));
     return true;
@@ -10362,7 +10366,8 @@ void CGOpenMPRuntime::registerTargetGlobalVariable(const VarDecl *VD,
   int64_t VarSize;
   llvm::GlobalValue::LinkageTypes Linkage;
 
-  if (*Res == OMPDeclareTargetDeclAttr::MT_To &&
+  if ((*Res == OMPDeclareTargetDeclAttr::MT_To ||
+       *Res == OMPDeclareTargetDeclAttr::MT_Enter) &&
       !HasRequiresUnifiedSharedMemory) {
     Flags = llvm::OffloadEntriesInfoManager::OMPTargetGlobalVarEntryTo;
     VarName = CGM.getMangledName(VD);
@@ -10393,7 +10398,8 @@ void CGOpenMPRuntime::registerTargetGlobalVariable(const VarDecl *VD,
     }
   } else {
     assert(((*Res == OMPDeclareTargetDeclAttr::MT_Link) ||
-            (*Res == OMPDeclareTargetDeclAttr::MT_To &&
+            ((*Res == OMPDeclareTargetDeclAttr::MT_To ||
+              *Res == OMPDeclareTargetDeclAttr::MT_Enter) &&
              HasRequiresUnifiedSharedMemory)) &&
            "Declare target attribute must link or to with unified memory.");
     if (*Res == OMPDeclareTargetDeclAttr::MT_Link)
@@ -10430,12 +10436,14 @@ void CGOpenMPRuntime::emitDeferredTargetDecls() const {
         OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD);
     if (!Res)
       continue;
-    if (*Res == OMPDeclareTargetDeclAttr::MT_To &&
+    if ((*Res == OMPDeclareTargetDeclAttr::MT_To ||
+         *Res == OMPDeclareTargetDeclAttr::MT_Enter) &&
         !HasRequiresUnifiedSharedMemory) {
       CGM.EmitGlobal(VD);
     } else {
       assert((*Res == OMPDeclareTargetDeclAttr::MT_Link ||
-              (*Res == OMPDeclareTargetDeclAttr::MT_To &&
+              ((*Res == OMPDeclareTargetDeclAttr::MT_To ||
+                *Res == OMPDeclareTargetDeclAttr::MT_Enter) &&
                HasRequiresUnifiedSharedMemory)) &&
              "Expected link clause or to clause with unified memory.");
       (void)CGM.getOpenMPRuntime().getAddrOfDeclareTargetVar(VD);
diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp
index efec43069ec2bc1..ae25767bc753aca 100644
--- a/clang/lib/CodeGen/CodeGenModule.cpp
+++ b/clang/lib/CodeGen/CodeGenModule.cpp
@@ -3305,12 +3305,14 @@ void CodeGenModule::EmitGlobal(GlobalDecl GD) {
                 OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD)) {
           bool UnifiedMemoryEnabled =
               getOpenMPRuntime().hasRequiresUnifiedSharedMemory();
-          if (*Res == OMPDeclareTargetDeclAttr::MT_To &&
+          if ((*Res == OMPDeclareTargetDeclAttr::MT_To ||
+               *Res == OMPDeclareTargetDeclAttr::MT_Enter) &&
               !UnifiedMemoryEnabled) {
             (void)GetAddrOfGlobalVar(VD);
           } else {
             assert(((*Res == OMPDeclareTargetDeclAttr::MT_Link) ||
-                    (*Res == OMPDeclareTargetDeclAttr::MT_To &&
+                    ((*Res == OMPDeclareTargetDeclAttr::MT_To ||
+                      *Res == OMPDeclareTargetDeclAttr::MT_Enter) &&
                      UnifiedMemoryEnabled)) &&
                    "Link clause or to clause with unified memory expected.");
             (void)getOpenMPRuntime().getAddrOfDeclareTargetVar(VD);
diff --git a/clang/lib/Parse/ParseOpenMP.cpp b/clang/lib/Parse/ParseOpenMP.cpp
index c7cdb348c38f9a5..82f86dae00843cb 100644
--- a/clang/lib/Parse/ParseOpenMP.cpp
+++ b/clang/lib/Parse/ParseOpenMP.cpp
@@ -1837,9 +1837,21 @@ void Parser::ParseOMPDeclareTargetClauses(
             << getOpenMPClauseName(OMPC_indirect) << 0;
         break;
       }
-      bool IsToOrLinkClause =
+      bool IsToEnterOrLinkClause =
           OMPDeclareTargetDeclAttr::ConvertStrToMapTypeTy(ClauseName, MT);
-      assert((!IsDeviceTypeClause || !IsToOrLinkClause) && "Cannot be both!");
+      assert((!IsDeviceTypeClause || !IsToEnterOrLinkClause) &&
+             "Cannot be both!");
+
+      // Starting with OpenMP 5.2 the `to` clause has been replaced by the
+      // `enter` clause.
+      if (getLangOpts().OpenMP >= 52 && ClauseName == "to") {
+        Diag(Tok, diag::err_omp_declare_target_unexpected_to_clause);
+        break;
+      }
+      if (getLangOpts().OpenMP <= 51 && ClauseName == "enter") {
+        Diag(Tok, diag::err_omp_declare_target_unexpected_enter_clause);
+        break;
+      }
 
       if (!IsDeviceTypeClause && !IsIndirectClause &&
           DTCI.Kind == OMPD_begin_declare_target) {
@@ -1847,16 +1859,18 @@ void Parser::ParseOMPDeclareTargetClauses(
             << ClauseName << (getLangOpts().OpenMP >= 51 ? 3 : 0);
         break;
       }
-      if (!IsDeviceTypeClause && !IsToOrLinkClause && !IsIndirectClause) {
-        Diag(Tok, diag::err_omp_declare_target_unexpected_clause)
+      if (!IsDeviceTypeClause && !IsToEnterOrLinkClause && !IsIndirectClause) {
+        Diag(Tok, getLangOpts().OpenMP >= 52
+                      ? diag::err_omp_declare_target_unexpected_clause_52
+                      : diag::err_omp_declare_target_unexpected_clause)
             << ClauseName
-            << (getLangOpts().OpenMP >= 51   ? 4
-                : getLangOpts().OpenMP >= 50 ? 2
-                                             : 1);
+            << (getLangOpts().OpenMP >= 51
+                    ? 4
+                    : getLangOpts().OpenMP >= 50 ? 2 : 1);
         break;
       }
 
-      if (IsToOrLinkClause || IsIndirectClause)
+      if (IsToEnterOrLinkClause || IsIndirectClause)
         HasToOrLinkOrIndirectClause = true;
 
       if (IsIndirectClause) {
@@ -1920,7 +1934,9 @@ void Parser::ParseOMPDeclareTargetClauses(
     }
     if (!HasIdentifier && Tok.isNot(tok::annot_pragma_openmp_end)) {
       Diag(Tok,
-           diag::err_omp_declare_target_unexpected_clause_after_implicit_to);
+           getLangOpts().OpenMP >= 52
+               ? diag::err_omp_declare_target_wrong_clause_after_implicit_enter
+               : diag::err_omp_declare_target_wrong_clause_after_implicit_to);
       break;
     }
 
@@ -1935,7 +1951,10 @@ void Parser::ParseOMPDeclareTargetClauses(
   // For declare target require at least 'to' or 'link' to be present.
   if (DTCI.Kind == OMPD_declare_target && RequiresToOrLinkOrIndirectClause &&
       !HasToOrLinkOrIndirectClause)
-    Diag(DTCI.Loc, diag::err_omp_declare_target_missing_to_or_link_clause)
+    Diag(DTCI.Loc,
+         getLangOpts().OpenMP >= 52
+             ? diag::err_omp_declare_target_missing_enter_or_link_clause
+             : diag::err_omp_declare_target_missing_to_or_link_clause)
         << (getLangOpts().OpenMP >= 51 ? 1 : 0);
 
   SkipUntil(tok::annot_pragma_openmp_end, StopBeforeMatch);
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index 8070f100608ffbf..26658a684d1d7f2 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -22710,7 +22710,8 @@ static void checkDeclInTargetContext(SourceLocation SL, SourceRange SR,
       (SemaRef.getCurLambda(/*IgnoreNonLambdaCapturingScope=*/true) ||
        SemaRef.getCurBlock() || SemaRef.getCurCapturedRegion()) &&
       VD->hasGlobalStorage()) {
-    if (!MapTy || *MapTy != OMPDeclareTargetDeclAttr::MT_To) {
+    if (!MapTy || (*MapTy != OMPDeclareTargetDeclAttr::MT_To &&
+                   *MapTy != OMPDeclareTargetDeclAttr::MT_Enter)) {
       // OpenMP 5.0, 2.12.7 declare target Directive, Restrictions
       // If a lambda declaration and definition appears between a
       // declare target directive and the matching end declare target
@@ -22791,8 +22792,11 @@ void Sema::checkDeclIsAllowedInOpenMPTarget(Expr *E, Decl *D,
             IsIndirect = true;
         }
         auto *A = OMPDeclareTargetDeclAttr::CreateImplicit(
-            Context, OMPDeclareTargetDeclAttr::MT_To, DTCI.DT, IndirectE,
-            IsIndirect, Level, SourceRange(DTCI.Loc, DTCI.Loc));
+            Context,
+            getLangOpts().OpenMP >= 52 ? OMPDeclareTargetDeclAttr::MT_Enter
+                                       : OMPDeclareTargetDeclAttr::MT_To,
+            DTCI.DT, IndirectE, IsIndirect, Level,
+            SourceRange(DTCI.Loc, DTCI.Loc));
         D->addAttr(A);
         if (ASTMutationListener *ML = Context.getASTMutationListener())
           ML->DeclarationMarkedOpenMPDeclareTarget(D, A);
diff --git a/clang/test/OpenMP/declare_target_ast_print.cpp b/clang/test/OpenMP/declare_target_ast_print.cpp
index 2dd45bbc79480d2..c0bc3445d0a7c07 100644
--- a/clang/test/OpenMP/declare_target_ast_print.cpp
+++ b/clang/test/OpenMP/declare_target_ast_print.cpp
@@ -4,10 +4,13 @@
 
 // RUN: %clang_cc1 -verify -fopenmp -I %S/Inputs -ast-print %s | FileCheck %s --check-prefix=CHECK --check-prefix=OMP50
 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=51 -I %S/Inputs -ast-print %s | FileCheck %s --check-prefix=CHECK --check-prefix=OMP51
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=52 -I %S/Inputs -ast-print %s | FileCheck %s --check-prefix=CHECK --check-prefix=OMP52
 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -I %S/Inputs -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp -std=c++11 -include-pch %t -fsyntax-only -I %S/Inputs -verify %s -ast-print | FileCheck %s --check-prefix=CHECK --check-prefix=OMP50
 // RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -I %S/Inputs -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -std=c++11 -include-pch %t -fsyntax-only -I %S/Inputs -verify %s -ast-print | FileCheck %s --check-prefix=CHECK --check-prefix=OMP51
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=52 -x c++ -std=c++11 -I %S/Inputs -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=52 -std=c++11 -include-pch %t -fsyntax-only -I %S/Inputs -verify %s -ast-print | FileCheck %s --check-prefix=CHECK --check-prefix=OMP52
 
 // RUN: %clang_cc1 -verify -fopenmp-simd -I %S/Inputs -ast-print %s | FileCheck %s
 // RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -I %S/Inputs -emit-pch -o %t %s
@@ -77,6 +80,48 @@ void xoo();
 }
 #endif // _OPENMP
 
+#if _OPENMP == 202111
+extern "C" {
+void boo_c() {}
+#pragma omp declare target enter(boo_c) indirect
+// OMP52: #pragma omp declare target indirect
+// OMP52: void boo_c() {
+// OMP52: }
+// OMP52: #pragma omp end declare target
+#pragma omp declare target indirect
+void yoo(){}
+#pragma omp end declare target
+// OMP52: #pragma omp declare target indirect
+// OMP52: void yoo() {
+// OMP52: }
+// OMP52: #pragma omp end declare target
+}
+extern "C++" {
+void boo_cpp() {}
+#pragma omp declare target enter(boo_cpp) indirect
+// OMP52: #pragma omp declare target indirect
+// OMP52: void boo_cpp() {
+// OMP52: }
+// OMP52: #pragma omp end declare target
+
+constexpr bool f() {return false;}
+#pragma omp begin declare target indirect(f())
+void zoo() {}
+void xoo();
+#pragma omp end declare target
+#pragma omp declare target enter(zoo) indirect(false)
+// OMP52: #pragma omp declare target indirect(f())
+// OMP52: #pragma omp declare target indirect(false)
+// OMP52: void zoo() {
+// OMP52: }
+// OMP52: #pragma omp end declare target
+// OMP52: #pragma omp declare target indirect(f())
+// OMP52: void xoo();
+// OMP52: #pragma omp end declare target
+
+}
+#endif // _OPENMP
+
 int out_decl_target = 0;
 #pragma omp declare target (out_decl_target)
 
@@ -185,7 +230,11 @@ void f1() {
 int b1, b2, b3;
 void f2() {
 }
+#if _OPENMP == 202111
+#pragma omp declare target enter(b1) enter(b2), enter(b3, f2)
+#else
 #pragma omp declare target to(b1) to(b2), to(b3, f2)
+#endif // _OPENMP == 202111
 // CHECK: #pragma omp declare target{{$}}
 // CHECK: int b1;
 // CHECK: #pragma omp end declare target{{$}}
@@ -288,7 +337,11 @@ int baz() { return 1; }
 
 #pragma omp declare target
 int abc1() { return 1; }
+#if _OPENMP == 202111
+#pragma omp declare target enter(abc1) device_type(nohost)
+#else
 #pragma omp declare target to(abc1) device_type(nohost)
+#endif // _OPENMP == 202111
 #pragma omp end declare target
 
 // CHECK-NEXT: #pragma omp declare target

>From 4a3eeea3f80aebdb4261a19f5c6bf4b8ab4dcfb9 Mon Sep 17 00:00:00 2001
From: Doru Bercea <doru.bercea at amd.com>
Date: Mon, 9 Oct 2023 20:52:20 -0400
Subject: [PATCH 2/2] Add support for map checks under unified shared memory

---
 openmp/libomptarget/include/device.h          |   4 +
 openmp/libomptarget/src/device.cpp            |  77 +++++-
 openmp/libomptarget/src/omptarget.cpp         | 156 ++++++------
 .../unified_map_checks_arrays.cpp             |  48 ++++
 .../unified_map_checks_close_enter_exit.cpp   | 223 ++++++++++++++++++
 .../unified_map_checks_close_modifier.cpp     | 184 +++++++++++++++
 .../unified_map_checks_error.cpp              |  36 +++
 .../unified_map_checks_no_target.cpp          |  36 +++
 .../unified_map_checks_scalars.cpp            |  99 ++++++++
 .../unified_map_checks_shared_update.cpp      | 137 +++++++++++
 .../zero_sized_array.cpp                      |  31 +++
 11 files changed, 958 insertions(+), 73 deletions(-)
 create mode 100644 openmp/libomptarget/test/unified_shared_memory/unified_map_checks_arrays.cpp
 create mode 100644 openmp/libomptarget/test/unified_shared_memory/unified_map_checks_close_enter_exit.cpp
 create mode 100644 openmp/libomptarget/test/unified_shared_memory/unified_map_checks_close_modifier.cpp
 create mode 100644 openmp/libomptarget/test/unified_shared_memory/unified_map_checks_error.cpp
 create mode 100644 openmp/libomptarget/test/unified_shared_memory/unified_map_checks_no_target.cpp
 create mode 100644 openmp/libomptarget/test/unified_shared_memory/unified_map_checks_scalars.cpp
 create mode 100644 openmp/libomptarget/test/unified_shared_memory/unified_map_checks_shared_update.cpp
 create mode 100644 openmp/libomptarget/test/unified_shared_memory/zero_sized_array.cpp

diff --git a/openmp/libomptarget/include/device.h b/openmp/libomptarget/include/device.h
index cd76d88618be4ee..56a4f5ba4242c18 100644
--- a/openmp/libomptarget/include/device.h
+++ b/openmp/libomptarget/include/device.h
@@ -447,6 +447,10 @@ struct DeviceTy {
   /// - Data allocation failed;
   /// - The user tried to do an illegal mapping;
   /// - Data transfer issue fails.
+  /// If unified shared memory is enabled the data will not be transferred to
+  /// the device and will be used from the host. Data will be added to the
+  /// mapping table to allow checks to happen even when in unified shared
+  /// memory.
   TargetPointerResultTy getTargetPointer(
       HDTTMapAccessorTy &HDTTMap, void *HstPtrBegin, void *HstPtrBase,
       int64_t TgtPadding, int64_t Size, map_var_info_t HstPtrName,
diff --git a/openmp/libomptarget/src/device.cpp b/openmp/libomptarget/src/device.cpp
index 93d2157dbd4ee15..a0946a733e66c48 100644
--- a/openmp/libomptarget/src/device.cpp
+++ b/openmp/libomptarget/src/device.cpp
@@ -268,6 +268,27 @@ TargetPointerResultTy DeviceTy::getTargetPointer(
          LR.TPR.getEntry()->holdRefCountToStr().c_str(), HoldRefCountAction,
          (HstPtrName) ? getNameFromMapping(HstPtrName).c_str() : "unknown");
     LR.TPR.TargetPointer = (void *)Ptr;
+
+    // When the target pointer is retrieved again, then the condition for this
+    // branch can be true hence preventing the unified shared memory to be
+    // taken at all. This ensures that the IsHostPointer and IsPresent flags
+    // are correctly set even in that situation.
+    if (PM->RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY &&
+        !HasCloseModifier && !LR.TPR.Flags.IsHostPointer) {
+      // This is a host pointer and is not present if the pointers match:
+      if (LR.TPR.getEntry()->TgtPtrBegin == LR.TPR.getEntry()->HstPtrBegin) {
+        LR.TPR.Flags.IsPresent = false;
+        LR.TPR.Flags.IsHostPointer = true;
+      }
+
+      // Catch the case where incoming HstPtrBegin is not consistent with the
+      // entry HstPtrBegin.
+      if (LR.TPR.Flags.IsHostPointer &&
+          ((uintptr_t)HstPtrBegin - LR.TPR.getEntry()->HstPtrBegin) != 0) {
+        assert(false &&
+               "Incoming HstPtrBegin different from entry HstPtrBegin");
+      }
+    }
   } else if ((LR.Flags.ExtendsBefore || LR.Flags.ExtendsAfter) && !IsImplicit) {
     // Explicit extension of mapped data - not allowed.
     MESSAGE("explicit extension not allowed: host address specified is " DPxMOD
@@ -289,13 +310,38 @@ TargetPointerResultTy DeviceTy::getTargetPointer(
     // In addition to the mapping rules above, the close map modifier forces the
     // mapping of the variable to the device.
     if (Size) {
-      DP("Return HstPtrBegin " DPxMOD " Size=%" PRId64 " for unified shared "
-         "memory\n",
-         DPxPTR((uintptr_t)HstPtrBegin), Size);
-      LR.TPR.Flags.IsPresent = false;
+      LR.TPR.Flags.IsNewEntry = true;
+      assert(TgtPadding == 0 && "TgtPadding must always be zero in USM mode");
+      uintptr_t TgtPtrBegin = (uintptr_t)HstPtrBegin + TgtPadding;
+      LR.TPR.setEntry(
+          HDTTMap
+              ->emplace(new HostDataToTargetTy(
+                  (uintptr_t)HstPtrBase, (uintptr_t)HstPtrBegin,
+                  (uintptr_t)HstPtrBegin + Size, (uintptr_t)HstPtrBegin,
+                  TgtPtrBegin, HasHoldModifier, HstPtrName))
+              .first->HDTT);
+      INFO(OMP_INFOTYPE_MAPPING_CHANGED, DeviceID,
+           "Creating new map entry ONLY with HstPtrBase=" DPxMOD
+           ", HstPtrBegin=" DPxMOD ", TgtAllocBegin=" DPxMOD
+           ", TgtPtrBegin=" DPxMOD
+           ", Size=%ld, DynRefCount=%s, HoldRefCount=%s, Name=%s\n",
+           DPxPTR(HstPtrBase), DPxPTR(HstPtrBegin), DPxPTR(HstPtrBegin),
+           DPxPTR(TgtPtrBegin), Size,
+           LR.TPR.getEntry()->dynRefCountToStr().c_str(),
+           LR.TPR.getEntry()->holdRefCountToStr().c_str(),
+           (HstPtrName) ? getNameFromMapping(HstPtrName).c_str() : "unknown");
       LR.TPR.Flags.IsHostPointer = true;
+
+      // The following assert should catch any case in which the pointers
+      // do not match to understand if this case can ever happen.
+      assert((uintptr_t)HstPtrBegin == TgtPtrBegin &&
+             "Pointers must always match");
+
+      // If the above assert is ever hit the following should be changed to =
+      // TgtPtrBegin
       LR.TPR.TargetPointer = HstPtrBegin;
     }
+    LR.TPR.Flags.IsPresent = false;
   } else if (HasPresentModifier) {
     DP("Mapping required by 'present' map type modifier does not exist for "
        "HstPtrBegin=" DPxMOD ", Size=%" PRId64 "\n",
@@ -444,6 +490,29 @@ DeviceTy::getTgtPtrBegin(void *HstPtrBegin, int64_t Size, bool UpdateRefCount,
          LR.TPR.getEntry()->dynRefCountToStr().c_str(), DynRefCountAction,
          LR.TPR.getEntry()->holdRefCountToStr().c_str(), HoldRefCountAction);
     LR.TPR.TargetPointer = (void *)TP;
+
+    // If this entry is not marked as being host pointer (the way the
+    // implementation works today this is never true, mistake?) then we
+    // have to check if this is a host pointer or not. This is a host pointer
+    // if the host address matches the target address.
+    if ((PM->RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY) &&
+        !LR.TPR.Flags.IsHostPointer) {
+      // If addresses match it means that we are dealing with a host pointer
+      // which has to be marked as one and present flag reset:
+      if (LR.TPR.getEntry()->TgtPtrBegin == LR.TPR.getEntry()->HstPtrBegin) {
+        LR.TPR.Flags.IsPresent = false;
+        LR.TPR.Flags.IsHostPointer = true;
+      }
+
+      // We want to catch the case where (uintptr_t)HstPtrBegin and
+      // LR.TPR.getEntry()->HstPtrBegin are not the same when LR is a host
+      // pointer. This case should never happen.
+      if (LR.TPR.Flags.IsHostPointer &&
+          ((uintptr_t)HstPtrBegin - LR.TPR.getEntry()->HstPtrBegin) != 0) {
+        assert(false &&
+               "Incoming HstPtrBegin different from entry HstPtrBegin");
+      }
+    }
   } else if (PM->RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY) {
     // If the value isn't found in the mapping and unified shared memory
     // is on then it means we have stumbled upon a value which we need to
diff --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp
index 40419e448942608..5a1fbdd031ffd45 100644
--- a/openmp/libomptarget/src/omptarget.cpp
+++ b/openmp/libomptarget/src/omptarget.cpp
@@ -788,7 +788,9 @@ postProcessingTargetDataEnd(DeviceTy *Device,
   int Ret = OFFLOAD_SUCCESS;
 
   for (auto &[HstPtrBegin, DataSize, ArgType, TPR] : EntriesInfo) {
-    bool DelEntry = !TPR.isHostPointer();
+    // Delete entry from the mapping table even when we are dealing with a
+    // host pointer.
+    bool DelEntry = true;
 
     // If the last element from the mapper (for end transfer args comes in
     // reverse order), do not remove the partial entry, the parent struct still
@@ -847,10 +849,12 @@ postProcessingTargetDataEnd(DeviceTy *Device,
     Ret = Device->eraseMapEntry(HDTTMap, Entry, DataSize);
     // Entry is already remove from the map, we can unlock it now.
     HDTTMap.destroy();
-    Ret |= Device->deallocTgtPtrAndEntry(Entry, DataSize);
-    if (Ret != OFFLOAD_SUCCESS) {
-      REPORT("Deallocating data from device failed.\n");
-      break;
+    if (!TPR.Flags.IsHostPointer) {
+      Ret |= Device->deallocTgtPtrAndEntry(Entry, DataSize);
+      if (Ret != OFFLOAD_SUCCESS) {
+        REPORT("Deallocating data from device failed.\n");
+        break;
+      }
     }
   }
 
@@ -909,78 +913,92 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
         Device.getTgtPtrBegin(HstPtrBegin, DataSize, UpdateRef, HasHoldModifier,
                               !IsImplicit, ForceDelete, /*FromDataEnd=*/true);
     void *TgtPtrBegin = TPR.TargetPointer;
-    if (!TPR.isPresent() && !TPR.isHostPointer() &&
-        (DataSize || HasPresentModifier)) {
-      DP("Mapping does not exist (%s)\n",
-         (HasPresentModifier ? "'present' map type modifier" : "ignored"));
-      if (HasPresentModifier) {
-        // OpenMP 5.1, sec. 2.21.7.1 "map Clause", p. 350 L10-13:
-        // "If a map clause appears on a target, target data, target enter data
-        // or target exit data construct with a present map-type-modifier then
-        // on entry to the region if the corresponding list item does not appear
-        // in the device data environment then an error occurs and the program
-        // terminates."
-        //
-        // This should be an error upon entering an "omp target exit data".  It
-        // should not be an error upon exiting an "omp target data" or "omp
-        // target".  For "omp target data", Clang thus doesn't include present
-        // modifiers for end calls.  For "omp target", we have not found a valid
-        // OpenMP program for which the error matters: it appears that, if a
-        // program can guarantee that data is present at the beginning of an
-        // "omp target" region so that there's no error there, that data is also
-        // guaranteed to be present at the end.
-        MESSAGE("device mapping required by 'present' map type modifier does "
-                "not exist for host address " DPxMOD " (%" PRId64 " bytes)",
-                DPxPTR(HstPtrBegin), DataSize);
-        return OFFLOAD_FAIL;
-      }
-    } else {
-      DP("There are %" PRId64 " bytes allocated at target address " DPxMOD
-         " - is%s last\n",
-         DataSize, DPxPTR(TgtPtrBegin), (TPR.Flags.IsLast ? "" : " not"));
-    }
-
-    // OpenMP 5.1, sec. 2.21.7.1 "map Clause", p. 351 L14-16:
-    // "If the map clause appears on a target, target data, or target exit data
-    // construct and a corresponding list item of the original list item is not
-    // present in the device data environment on exit from the region then the
-    // list item is ignored."
-    if (!TPR.isPresent())
-      continue;
 
-    // Move data back to the host
-    const bool HasAlways = ArgTypes[I] & OMP_TGT_MAPTYPE_ALWAYS;
-    const bool HasFrom = ArgTypes[I] & OMP_TGT_MAPTYPE_FROM;
-    if (HasFrom && (HasAlways || TPR.Flags.IsLast) &&
-        !TPR.Flags.IsHostPointer && DataSize != 0) {
-      DP("Moving %" PRId64 " bytes (tgt:" DPxMOD ") -> (hst:" DPxMOD ")\n",
-         DataSize, DPxPTR(TgtPtrBegin), DPxPTR(HstPtrBegin));
-
-      // Wait for any previous transfer if an event is present.
-      if (void *Event = TPR.getEntry()->getEvent()) {
-        if (Device.waitEvent(Event, AsyncInfo) != OFFLOAD_SUCCESS) {
-          REPORT("Failed to wait for event " DPxMOD ".\n", DPxPTR(Event));
+    // Check if HstPtrBegin matches the State HstPtrBegin or if any HstPtrBegin
+    // values have been registered:
+    bool HostPointerMismatch = true;
+    if (TPR.getEntry())
+      HostPointerMismatch =
+          TPR.getEntry()->HstPtrBegin != (uintptr_t)HstPtrBegin;
+
+    if (!TPR.isHostPointer()) {
+      if (!TPR.isPresent() && (DataSize || HasPresentModifier)) {
+        DP("Mapping does not exist (%s)\n",
+           (HasPresentModifier ? "'present' map type modifier" : "ignored"));
+        if (HasPresentModifier) {
+          // OpenMP 5.1, sec. 2.21.7.1 "map Clause", p. 350 L10-13:
+          // "If a map clause appears on a target, target data, target enter
+          // data or target exit data construct with a present map-type-modifier
+          // then on entry to the region if the corresponding list item does not
+          // appear in the device data environment then an error occurs and the
+          // program terminates."
+          //
+          // This should be an error upon entering an "omp target exit data". It
+          // should not be an error upon exiting an "omp target data" or "omp
+          // target".  For "omp target data", Clang thus doesn't include present
+          // modifiers for end calls.  For "omp target", we have not found a
+          // valid OpenMP program for which the error matters: it appears that,
+          // if a program can guarantee that data is present at the beginning of
+          // an "omp target" region so that there's no error there, that data is
+          // also guaranteed to be present at the end.
+          MESSAGE("device mapping required by 'present' map type modifier does "
+                  "not exist for host address " DPxMOD " (%" PRId64 " bytes)",
+                  DPxPTR(HstPtrBegin), DataSize);
           return OFFLOAD_FAIL;
         }
+      } else {
+        DP("There are %" PRId64 " bytes allocated at target address " DPxMOD
+           " - is%s last\n",
+           DataSize, DPxPTR(TgtPtrBegin), (TPR.Flags.IsLast ? "" : " not"));
       }
 
-      Ret = Device.retrieveData(HstPtrBegin, TgtPtrBegin, DataSize, AsyncInfo,
-                                TPR.getEntry());
-      if (Ret != OFFLOAD_SUCCESS) {
-        REPORT("Copying data from device failed.\n");
-        return OFFLOAD_FAIL;
-      }
+      // OpenMP 5.1, sec. 2.21.7.1 "map Clause", p. 351 L14-16:
+      // "If the map clause appears on a target, target data, or target exit
+      // data construct and a corresponding list item of the original list item
+      // is not present in the device data environment on exit from the region
+      // then the list item is ignored."
+      if (!TPR.isPresent())
+        continue;
 
-      // As we are expecting to delete the entry the d2h copy might race
-      // with another one that also tries to delete the entry. This happens
-      // as the entry can be reused and the reuse might happen after the
-      // copy-back was issued but before it completed. Since the reuse might
-      // also copy-back a value we would race.
-      if (TPR.Flags.IsLast) {
-        if (TPR.getEntry()->addEventIfNecessary(Device, AsyncInfo) !=
-            OFFLOAD_SUCCESS)
+      // Move data back to the host
+      const bool HasAlways = ArgTypes[I] & OMP_TGT_MAPTYPE_ALWAYS;
+      const bool HasFrom = ArgTypes[I] & OMP_TGT_MAPTYPE_FROM;
+      if (HasFrom && (HasAlways || TPR.Flags.IsLast) && DataSize != 0) {
+        DP("Moving %" PRId64 " bytes (tgt:" DPxMOD ") -> (hst:" DPxMOD ")\n",
+           DataSize, DPxPTR(TgtPtrBegin), DPxPTR(HstPtrBegin));
+
+        // Wait for any previous transfer if an event is present.
+        if (void *Event = TPR.getEntry()->getEvent()) {
+          if (Device.waitEvent(Event, AsyncInfo) != OFFLOAD_SUCCESS) {
+            REPORT("Failed to wait for event " DPxMOD ".\n", DPxPTR(Event));
+            return OFFLOAD_FAIL;
+          }
+        }
+
+        Ret = Device.retrieveData(HstPtrBegin, TgtPtrBegin, DataSize, AsyncInfo,
+                                  TPR.getEntry());
+        if (Ret != OFFLOAD_SUCCESS) {
+          REPORT("Copying data from device failed.\n");
           return OFFLOAD_FAIL;
+        }
+
+        // As we are expecting to delete the entry the d2h copy might race
+        // with another one that also tries to delete the entry. This happens
+        // as the entry can be reused and the reuse might happen after the
+        // copy-back was issued but before it completed. Since the reuse might
+        // also copy-back a value we would race.
+        if (TPR.Flags.IsLast) {
+          if (TPR.getEntry()->addEventIfNecessary(Device, AsyncInfo) !=
+              OFFLOAD_SUCCESS)
+            return OFFLOAD_FAIL;
+        }
       }
+    } else {
+      // Some zero-sized arrays are not mapped or added to the mapping table so
+      // they do not need to be removed. These arrays are not part of the
+      // current entry.
+      if (DataSize == 0 && !TPR.isPresent() && HostPointerMismatch)
+        continue;
     }
 
     // Add pointer to the buffer for post-synchronize processing.
diff --git a/openmp/libomptarget/test/unified_shared_memory/unified_map_checks_arrays.cpp b/openmp/libomptarget/test/unified_shared_memory/unified_map_checks_arrays.cpp
new file mode 100644
index 000000000000000..1b1d309b6939fcc
--- /dev/null
+++ b/openmp/libomptarget/test/unified_shared_memory/unified_map_checks_arrays.cpp
@@ -0,0 +1,48 @@
+// clang-format off
+// RUN: %libomptarget-compilexx-generic && env HSA_XNACK=1 LIBOMPTARGET_INFO=-1 %libomptarget-run-generic 2>&1 | %fcheck-generic
+// clang-format on
+
+// REQUIRES: amdgcn-amd-amdhsa
+// UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
+
+#include <omp.h>
+
+#include <cassert>
+#include <iostream>
+
+#pragma omp requires unified_shared_memory
+
+int main(int argc, char *argv[]) {
+  int *v = (int *)malloc(sizeof(int) * 10);
+
+// clang-format off
+// CHECK: Entering OpenMP kernel at {{.*}} with 1 arguments:
+// CHECK: Creating new map entry ONLY with HstPtrBase=[[V_HST_PTR_ADDR:0x.*]], HstPtrBegin=[[V_HST_PTR_ADDR]], TgtAllocBegin=[[V_HST_PTR_ADDR]], TgtPtrBegin=[[V_HST_PTR_ADDR]], Size=200, DynRefCount=1, HoldRefCount=0
+// CHECK: Mapping exists with HstPtrBegin=[[V_HST_PTR_ADDR]], TgtPtrBegin=[[V_HST_PTR_ADDR]], Size=200, DynRefCount=1 (update suppressed), HoldRefCount=0
+// CHECK: Launching kernel __omp_offloading_{{.*}}_main_l{{.*}} with 1 blocks and 256 threads in Generic mode
+
+// CHECK: Mapping exists with HstPtrBegin=[[V_HST_PTR_ADDR]], TgtPtrBegin=[[V_HST_PTR_ADDR]], Size=200, DynRefCount=0 (decremented, delayed deletion), HoldRefCount=0
+// CHECK: Removing map entry with HstPtrBegin=[[V_HST_PTR_ADDR]], TgtPtrBegin=[[V_HST_PTR_ADDR]], Size=200
+
+// CHECK: Entering OpenMP kernel at {{.*}} with 1 arguments:
+// CHECK: Creating new map entry ONLY with HstPtrBase=[[V_HST_PTR_ADDR]], HstPtrBegin=[[V_HST_PTR_ADDR]], TgtAllocBegin=[[V_HST_PTR_ADDR]], TgtPtrBegin=[[V_HST_PTR_ADDR]], Size=280, DynRefCount=1, HoldRefCount=0
+// CHECK: Mapping exists with HstPtrBegin=[[V_HST_PTR_ADDR]], TgtPtrBegin=[[V_HST_PTR_ADDR]], Size=280, DynRefCount=1 (update suppressed), HoldRefCount=0
+// CHECK: Launching kernel __omp_offloading_{{.*}}_main_l{{.*}} with 1 blocks and 256 threads in Generic mode
+// CHECK: Mapping exists with HstPtrBegin=[[V_HST_PTR_ADDR]], TgtPtrBegin=[[V_HST_PTR_ADDR]], Size=280, DynRefCount=0 (decremented, delayed deletion), HoldRefCount=0
+// CHECK: Removing map entry with HstPtrBegin=[[V_HST_PTR_ADDR]], TgtPtrBegin=[[V_HST_PTR_ADDR]], Size=280
+// clang-format on
+#pragma omp target map(tofrom : v[ : 50])
+  { v[32] = 32; }
+
+#pragma omp target map(tofrom : v[ : 70])
+  { v[64] = 64; }
+
+  printf("v[32] = %d, v[64] = %d\n", v[32], v[64]);
+
+  free(v);
+
+  std::cout << "PASS\n";
+  return 0;
+}
+// CHECK: v[32] = 32, v[64] = 64
+// CHECK: PASS
diff --git a/openmp/libomptarget/test/unified_shared_memory/unified_map_checks_close_enter_exit.cpp b/openmp/libomptarget/test/unified_shared_memory/unified_map_checks_close_enter_exit.cpp
new file mode 100644
index 000000000000000..b20f997caf209e2
--- /dev/null
+++ b/openmp/libomptarget/test/unified_shared_memory/unified_map_checks_close_enter_exit.cpp
@@ -0,0 +1,223 @@
+// clang-format off
+// RUN: %libomptarget-compilexx-generic && env HSA_XNACK=1 LIBOMPTARGET_INFO=-1 %libomptarget-run-generic 2>&1 | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
+
+// REQUIRES: amdgcn-amd-amdhsa
+
+#include <omp.h>
+#include <stdio.h>
+
+#pragma omp requires unified_shared_memory
+
+#define N 1024
+
+int main(int argc, char *argv[]) {
+  int fails;
+  void *host_alloc = 0, *device_alloc = 0;
+  int *a = (int *)malloc(N * sizeof(int));
+  int dev = omp_get_default_device();
+
+  // Init
+  for (int i = 0; i < N; ++i) {
+    a[i] = 10;
+  }
+  host_alloc = &a[0];
+
+  //
+  // map + target no close
+  //
+
+// clang-format off
+// CHECK: Entering OpenMP data region with being_mapper at {{.*}} with 2 arguments:
+// CHECK: Creating new map entry ONLY with HstPtrBase=[[A_HST_PTR:0x.*]], HstPtrBegin=[[A_HST_PTR]], TgtAllocBegin=[[A_HST_PTR]], TgtPtrBegin=[[A_HST_PTR]], Size=4096, DynRefCount=1, HoldRefCount=0
+// CHECK: Creating new map entry ONLY with HstPtrBase=[[DEVICE_ALLOC_HST_PTR:0x.*]], HstPtrBegin=[[DEVICE_ALLOC_HST_PTR]], TgtAllocBegin=[[DEVICE_ALLOC_HST_PTR]], TgtPtrBegin=[[DEVICE_ALLOC_HST_PTR]], Size=8, DynRefCount=1, HoldRefCount=0
+// CHECK: OpenMP Host-Device pointer mappings after block
+// CHECK: Host Ptr
+// CHECK: [[A_HST_PTR]]
+// CHECK: [[DEVICE_ALLOC_HST_PTR]]
+// clang-format on
+#pragma omp target data map(tofrom : a[ : N]) map(tofrom : device_alloc)
+  {
+// clang-format off
+// CHECK: Entering OpenMP kernel at {{.*}} with 2 arguments:
+// CHECK: Mapping exists with HstPtrBegin=[[DEVICE_ALLOC_HST_PTR]], TgtPtrBegin=[[DEVICE_ALLOC_HST_PTR]], Size=8, DynRefCount=2 (incremented), HoldRefCount=0
+// CHECK: Mapping exists (implicit) with HstPtrBegin=[[A_HST_PTR]], TgtPtrBegin=[[A_HST_PTR]], Size=0, DynRefCount=2 (incremented), HoldRefCount=0
+// CHECK: Mapping exists with HstPtrBegin=[[DEVICE_ALLOC_HST_PTR]], TgtPtrBegin=[[DEVICE_ALLOC_HST_PTR]], Size=8, DynRefCount=2 (update suppressed), HoldRefCount=0
+// CHECK: Mapping exists with HstPtrBegin=[[A_HST_PTR]], TgtPtrBegin=[[A_HST_PTR]], Size=0, DynRefCount=2 (update suppressed), HoldRefCount=0
+// CHECK: Launching kernel __omp_offloading_{{.*}}_main_l{{.*}} with 1 blocks and 256 threads in Generic mode
+// clang-format on
+#pragma omp target map(tofrom : device_alloc)
+    { device_alloc = &a[0]; }
+  }
+// clang-format off
+// CHECK: Mapping exists with HstPtrBegin=[[A_HST_PTR]], TgtPtrBegin=[[A_HST_PTR]], Size=0, DynRefCount=1 (decremented), HoldRefCount=0
+// CHECK: Mapping exists with HstPtrBegin=[[DEVICE_ALLOC_HST_PTR]], TgtPtrBegin=[[DEVICE_ALLOC_HST_PTR]], Size=8, DynRefCount=1 (decremented), HoldRefCount=0
+// CHECK: OpenMP Host-Device pointer mappings after block
+// CHECK: Host Ptr
+// CHECK: [[A_HST_PTR]]
+// CHECK: [[DEVICE_ALLOC_HST_PTR]]
+// CHECK: Exiting OpenMP data region with end_mapper at {{.*}} with 2 arguments:
+// CHECK: Mapping exists with HstPtrBegin=[[DEVICE_ALLOC_HST_PTR]], TgtPtrBegin=[[DEVICE_ALLOC_HST_PTR]], Size=8, DynRefCount=0 (decremented, delayed deletion), HoldRefCount=0
+// CHECK: Mapping exists with HstPtrBegin=[[A_HST_PTR]], TgtPtrBegin=[[A_HST_PTR]], Size=4096, DynRefCount=0 (decremented, delayed deletion), HoldRefCount=0
+// CHECK: Removing map entry with HstPtrBegin=[[DEVICE_ALLOC_HST_PTR]], TgtPtrBegin=[[DEVICE_ALLOC_HST_PTR]], Size=8
+// CHECK: Removing map entry with HstPtrBegin=[[A_HST_PTR]], TgtPtrBegin=[[A_HST_PTR]], Size=4096
+// clang-format on
+  if (device_alloc == host_alloc)
+    printf("a used from unified memory.\n");
+
+  //
+  // map + target with close
+  //
+// clang-format off
+// CHECK: Entering OpenMP data region with being_mapper at {{.*}} with 2 arguments:
+// CHECK: Creating new map entry with HstPtrBase=[[A_HST_PTR]], HstPtrBegin=[[A_HST_PTR]], TgtAllocBegin=[[A_DEV_PTR:0x.*]], TgtPtrBegin=[[A_DEV_PTR]], Size=4096, DynRefCount=1, HoldRefCount=0
+// CHECK: Copying data from host to device, HstPtr=[[A_HST_PTR]], TgtPtr=[[A_DEV_PTR]], Size=4096
+// CHECK: Creating new map entry ONLY with HstPtrBase=[[DEVICE_ALLOC_HST_PTR]], HstPtrBegin=[[DEVICE_ALLOC_HST_PTR]], TgtAllocBegin=[[DEVICE_ALLOC_HST_PTR]], TgtPtrBegin=[[DEVICE_ALLOC_HST_PTR]], Size=8, DynRefCount=1, HoldRefCount=0
+// CHECK: OpenMP Host-Device pointer mappings after block
+// CHECK: Host Ptr
+// CHECK: [[A_HST_PTR]]
+// CHECK: [[DEVICE_ALLOC_HST_PTR]]
+// clang-format on
+  device_alloc = 0;
+#pragma omp target data map(close, tofrom : a[ : N]) map(tofrom : device_alloc)
+  {
+// clang-format off
+// CHECK: Entering OpenMP kernel at {{.*}} with 2 arguments:
+// CHECK: Mapping exists with HstPtrBegin=[[DEVICE_ALLOC_HST_PTR]], TgtPtrBegin=[[DEVICE_ALLOC_HST_PTR]], Size=8, DynRefCount=2 (incremented), HoldRefCount=0
+// CHECK: Mapping exists (implicit) with HstPtrBegin=[[A_HST_PTR]], TgtPtrBegin=[[A_DEV_PTR]], Size=0, DynRefCount=2 (incremented), HoldRefCount=0
+// CHECK: Mapping exists with HstPtrBegin=[[DEVICE_ALLOC_HST_PTR]], TgtPtrBegin=[[DEVICE_ALLOC_HST_PTR]], Size=8, DynRefCount=2 (update suppressed), HoldRefCount=0
+// CHECK: Mapping exists with HstPtrBegin=[[A_HST_PTR]], TgtPtrBegin=[[A_DEV_PTR]], Size=0, DynRefCount=2 (update suppressed), HoldRefCount=0
+// CHECK: Launching kernel __omp_offloading_{{.*}}_main_l{{.*}} with 1 blocks and 256 threads in Generic mode
+// clang-format on
+#pragma omp target map(tofrom : device_alloc)
+    { device_alloc = &a[0]; }
+  }
+// clang-format off
+// CHECK: Mapping exists with HstPtrBegin=[[A_HST_PTR]], TgtPtrBegin=[[A_DEV_PTR]], Size=0, DynRefCount=1 (decremented), HoldRefCount=0
+// CHECK: Mapping exists with HstPtrBegin=[[DEVICE_ALLOC_HST_PTR]], TgtPtrBegin=[[DEVICE_ALLOC_HST_PTR]], Size=8, DynRefCount=1 (decremented), HoldRefCount=0
+// CHECK: OpenMP Host-Device pointer mappings after block
+// CHECK: Host Ptr
+// CHECK: [[A_HST_PTR]]
+// CHECK: [[DEVICE_ALLOC_HST_PTR]]
+// CHECK: Exiting OpenMP data region with end_mapper at {{.*}} with 2 arguments:
+// CHECK: Mapping exists with HstPtrBegin=[[DEVICE_ALLOC_HST_PTR]], TgtPtrBegin=[[DEVICE_ALLOC_HST_PTR]], Size=8, DynRefCount=0 (decremented, delayed deletion), HoldRefCount=0
+// CHECK: Mapping exists with HstPtrBegin=[[A_HST_PTR]], TgtPtrBegin=[[A_DEV_PTR]], Size=4096, DynRefCount=0 (decremented, delayed deletion), HoldRefCount=0
+// CHECK: Copying data from device to host, TgtPtr=[[A_DEV_PTR]], HstPtr=[[A_HST_PTR]], Size=4096
+// CHECK: Removing map entry with HstPtrBegin=[[DEVICE_ALLOC_HST_PTR]], TgtPtrBegin=[[DEVICE_ALLOC_HST_PTR]], Size=8
+// CHECK: Removing map entry with HstPtrBegin=[[A_HST_PTR]], TgtPtrBegin=[[A_DEV_PTR]], Size=4096
+// clang-format on
+  if (device_alloc != host_alloc)
+    printf("a copied to device.\n");
+
+  //
+  // map + use_device_ptr no close
+  //
+// clang-format off
+// CHECK: Entering OpenMP data region with being_mapper at {{.*}} with 1 arguments:
+// CHECK: Creating new map entry ONLY with HstPtrBase=[[A_HST_PTR]], HstPtrBegin=[[A_HST_PTR]], TgtAllocBegin=[[A_HST_PTR]], TgtPtrBegin=[[A_HST_PTR]], Size=4096, DynRefCount=1, HoldRefCount=0
+// CHECK: OpenMP Host-Device pointer mappings after block
+// CHECK: Host Ptr
+// CHECK: [[A_HST_PTR]]
+// clang-format on
+  device_alloc = 0;
+#pragma omp target data map(tofrom : a[ : N]) use_device_ptr(a)
+  { device_alloc = &a[0]; }
+// clang-format off
+// CHECK: Exiting OpenMP data region with end_mapper at {{.*}} with 1 arguments:
+// CHECK: Mapping exists with HstPtrBegin=[[A_HST_PTR]], TgtPtrBegin=[[A_HST_PTR]], Size=4096, DynRefCount=0 (decremented, delayed deletion), HoldRefCount=0
+// CHECK: Removing map entry with HstPtrBegin=[[A_HST_PTR]], TgtPtrBegin=[[A_HST_PTR]], Size=4096
+// clang-format on
+  if (device_alloc == host_alloc)
+    printf("a used from unified memory with use_device_ptr.\n");
+
+  //
+  // map + use_device_ptr close
+  //
+// clang-format off
+// CHECK: Entering OpenMP data region with being_mapper at {{.*}} with 1 arguments:
+// CHECK: Creating new map entry with HstPtrBase=[[A_HST_PTR]], HstPtrBegin=[[A_HST_PTR]], TgtAllocBegin=[[A_DEV_PTR:0x.*]], TgtPtrBegin=[[A_DEV_PTR]], Size=4096, DynRefCount=1, HoldRefCount=0
+// CHECK: Copying data from host to device, HstPtr=[[A_HST_PTR]], TgtPtr=[[A_DEV_PTR]], Size=4096
+// CHECK: OpenMP Host-Device pointer mappings after block
+// CHECK: Host Ptr
+// CHECK: [[A_HST_PTR]]
+// clang-format on
+  device_alloc = 0;
+#pragma omp target data map(close, tofrom : a[ : N]) use_device_ptr(a)
+  { device_alloc = &a[0]; }
+// clang-format off
+// CHECK: Exiting OpenMP data region with end_mapper at {{.*}} with 1 arguments:
+// CHECK: Mapping exists with HstPtrBegin=[[A_HST_PTR]], TgtPtrBegin=[[A_DEV_PTR]], Size=4096, DynRefCount=0 (decremented, delayed deletion), HoldRefCount=0
+// CHECK: Copying data from device to host, TgtPtr=[[A_DEV_PTR]], HstPtr=[[A_HST_PTR]], Size=4096
+// CHECK: Removing map entry with HstPtrBegin=[[A_HST_PTR]], TgtPtrBegin=[[A_DEV_PTR]], Size=4096
+// clang-format on
+  if (device_alloc != host_alloc)
+    printf("a used from device memory with use_device_ptr.\n");
+
+  //
+  // map enter/exit + close
+  //
+// clang-format off
+// CHECK: Entering OpenMP data region with being_mapper at {{.*}} with 1 arguments:
+// CHECK: Creating new map entry with HstPtrBase=[[A_HST_PTR]], HstPtrBegin=[[A_HST_PTR]], TgtAllocBegin=[[A_DEV_PTR:0x.*]], TgtPtrBegin=[[A_DEV_PTR]], Size=4096, DynRefCount=1, HoldRefCount=0
+// CHECK: Copying data from host to device, HstPtr=[[A_HST_PTR]], TgtPtr=[[A_DEV_PTR]], Size=4096
+// CHECK: OpenMP Host-Device pointer mappings after block
+// CHECK: Host Ptr
+// CHECK: [[A_HST_PTR]]
+// clang-format on
+  device_alloc = 0;
+#pragma omp target enter data map(close, to : a[ : N])
+// clang-format off
+// CHECK: Entering OpenMP kernel at {{.*}} with 2 arguments:
+// CHECK: Creating new map entry ONLY with HstPtrBase=[[DEVICE_ALLOC_HST_PTR]], HstPtrBegin=[[DEVICE_ALLOC_HST_PTR]], TgtAllocBegin=[[DEVICE_ALLOC_HST_PTR]], TgtPtrBegin=[[DEVICE_ALLOC_HST_PTR]], Size=8, DynRefCount=1, HoldRefCount=0
+// CHECK: Mapping exists (implicit) with HstPtrBegin=[[A_HST_PTR]], TgtPtrBegin=[[A_DEV_PTR]], Size=0, DynRefCount=2 (incremented), HoldRefCount=0
+// CHECK: Mapping exists with HstPtrBegin=[[DEVICE_ALLOC_HST_PTR]], TgtPtrBegin=[[DEVICE_ALLOC_HST_PTR]], Size=8, DynRefCount=1 (update suppressed), HoldRefCount=0
+// CHECK: Mapping exists with HstPtrBegin=[[A_HST_PTR]], TgtPtrBegin=[[A_DEV_PTR]], Size=0, DynRefCount=2 (update suppressed), HoldRefCount=0
+// CHECK: Launching kernel __omp_offloading_{{.*}}_main_l{{.*}} with 1 blocks and 256 threads in Generic mode
+// clang-format on
+#pragma omp target map(from : device_alloc)
+  {
+    device_alloc = &a[0];
+    a[0] = 99;
+  }
+// clang-format off
+// CHECK: Mapping exists with HstPtrBegin=[[A_HST_PTR]], TgtPtrBegin=[[A_DEV_PTR]], Size=0, DynRefCount=1 (decremented), HoldRefCount=0
+// CHECK: Mapping exists with HstPtrBegin=[[DEVICE_ALLOC_HST_PTR]], TgtPtrBegin=[[DEVICE_ALLOC_HST_PTR]], Size=8, DynRefCount=0 (decremented, delayed deletion), HoldRefCount=0
+// CHECK: Removing map entry with HstPtrBegin=[[DEVICE_ALLOC_HST_PTR]], TgtPtrBegin=[[DEVICE_ALLOC_HST_PTR]], Size=8
+// CHECK: OpenMP Host-Device pointer mappings after block
+// CHECK: Host Ptr
+// CHECK: [[A_HST_PTR]]
+// CHECK: Exiting OpenMP data region with end_mapper at {{.*}} with 1 arguments:
+// CHECK: Mapping exists with HstPtrBegin=[[A_HST_PTR]], TgtPtrBegin=[[A_DEV_PTR]], Size=4096, DynRefCount=0 (decremented, delayed deletion), HoldRefCount=0
+// CHECK: Copying data from device to host, TgtPtr=[[A_DEV_PTR]], HstPtr=[[A_HST_PTR]], Size=4096
+// CHECK: Removing map entry with HstPtrBegin=[[A_HST_PTR]], TgtPtrBegin=[[A_DEV_PTR]], Size=4096
+// clang-format on
+
+  // 'close' is missing, so the runtime must check whether s is actually in
+  // shared memory in order to determine whether to transfer data and delete the
+  // allocation.
+#pragma omp target exit data map(from : a[ : N])
+
+  if (device_alloc != host_alloc)
+    printf("a has been mapped to the device.\n");
+
+  printf("a[0]=%d\n", a[0]);
+  printf("a is present: %d\n", omp_target_is_present(a, dev));
+
+  free(a);
+
+  // CHECK: a used from unified memory.
+  // CHECK: a copied to device.
+  // CHECK: a used from unified memory with use_device_ptr.
+
+  // CHECK: a used from device memory with use_device_ptr.
+  // CHECK: a has been mapped to the device.
+  // CHECK: a[0]=99
+  // CHECK: a is present: 0
+
+  // CHECK: Done!
+  printf("Done!\n");
+
+  return 0;
+}
diff --git a/openmp/libomptarget/test/unified_shared_memory/unified_map_checks_close_modifier.cpp b/openmp/libomptarget/test/unified_shared_memory/unified_map_checks_close_modifier.cpp
new file mode 100644
index 000000000000000..2969d42c9389b44
--- /dev/null
+++ b/openmp/libomptarget/test/unified_shared_memory/unified_map_checks_close_modifier.cpp
@@ -0,0 +1,184 @@
+// clang-format off
+// RUN: %libomptarget-compilexx-generic && env HSA_XNACK=1 LIBOMPTARGET_INFO=-1 %libomptarget-run-generic 2>&1 | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
+
+// REQUIRES: amdgcn-amd-amdhsa
+
+#include <omp.h>
+#include <stdio.h>
+
+#pragma omp requires unified_shared_memory
+
+#define N 1024
+
+int main(int argc, char *argv[]) {
+  int fails;
+  void *host_alloc, *device_alloc;
+  void *host_data, *device_data;
+  int *alloc = (int *)malloc(N * sizeof(int));
+  int data[N];
+
+  for (int i = 0; i < N; ++i) {
+    alloc[i] = 10;
+    data[i] = 1;
+  }
+
+  host_data = &data[0];
+  host_alloc = &alloc[0];
+
+  //
+  // Test that updates on the device are not visible to host
+  // when only a TO mapping is used.
+  //
+
+// clang-format off
+// CHECK: Creating new map entry ONLY with HstPtrBase=[[DEVICE_DATA_HST_PTR:0x.*]], HstPtrBegin=[[DEVICE_DATA_HST_PTR]], TgtAllocBegin=[[DEVICE_DATA_HST_PTR]], TgtPtrBegin=[[DEVICE_DATA_HST_PTR]], Size=8, DynRefCount=1, HoldRefCount=0
+// CHECK: Creating new map entry with HstPtrBase=[[DATA_HST_PTR:0x.*]], HstPtrBegin=[[DATA_HST_PTR]], TgtAllocBegin=[[DATA_DEV_PTR:0x.*]], TgtPtrBegin=[[DATA_DEV_PTR]], Size=4096, DynRefCount=1, HoldRefCount=0
+// CHECK: Copying data from host to device, HstPtr=[[DATA_HST_PTR]], TgtPtr=[[DATA_DEV_PTR]], Size=4096
+// CHECK: Creating new map entry ONLY with HstPtrBase=[[DEVICE_ALLOC_HST_PTR:0x.*]], HstPtrBegin=[[DEVICE_ALLOC_HST_PTR]], TgtAllocBegin=[[DEVICE_ALLOC_HST_PTR]], TgtPtrBegin=[[DEVICE_ALLOC_HST_PTR]], Size=8, DynRefCount=1, HoldRefCount=0
+// CHECK: Creating new map entry with HstPtrBase=[[ALLOC_HST_PTR:0x.*]], HstPtrBegin=[[ALLOC_HST_PTR]], TgtAllocBegin=[[ALLOC_DEV_PTR:0x.*]], TgtPtrBegin=[[ALLOC_DEV_PTR]], Size=4096, DynRefCount=1, HoldRefCount=0
+// CHECK: Copying data from host to device, HstPtr=[[ALLOC_HST_PTR]], TgtPtr=[[ALLOC_DEV_PTR]], Size=4096
+
+// CHECK: Mapping exists with HstPtrBegin=[[DEVICE_DATA_HST_PTR]], TgtPtrBegin=[[DEVICE_DATA_HST_PTR]], Size=8, DynRefCount=1 (update suppressed), HoldRefCount=0
+// CHECK: Mapping exists with HstPtrBegin=[[DATA_HST_PTR]], TgtPtrBegin=[[DATA_DEV_PTR]], Size=4096, DynRefCount=1 (update suppressed), HoldRefCount=0
+// CHECK: Mapping exists with HstPtrBegin=[[DEVICE_ALLOC_HST_PTR]], TgtPtrBegin=[[DEVICE_ALLOC_HST_PTR]], Size=8, DynRefCount=1 (update suppressed), HoldRefCount=0
+// CHECK: Mapping exists with HstPtrBegin=[[ALLOC_HST_PTR]], TgtPtrBegin=[[ALLOC_DEV_PTR]], Size=4096, DynRefCount=1 (update suppressed), HoldRefCount=0
+// CHECK: Launching kernel __omp_offloading_{{.*}}_main_l{{.*}} with 1 blocks and 256 threads in Generic mode
+
+// CHECK: Mapping exists with HstPtrBegin=[[ALLOC_HST_PTR]], TgtPtrBegin=[[ALLOC_DEV_PTR]], Size=4096, DynRefCount=0 (decremented, delayed deletion), HoldRefCount=0
+// CHECK: Mapping exists with HstPtrBegin=[[DEVICE_ALLOC_HST_PTR]], TgtPtrBegin=[[DEVICE_ALLOC_HST_PTR]], Size=8, DynRefCount=0 (decremented, delayed deletion), HoldRefCount=0
+// CHECK: Mapping exists with HstPtrBegin=[[DATA_HST_PTR]], TgtPtrBegin=[[DATA_DEV_PTR]], Size=4096, DynRefCount=0 (decremented, delayed deletion), HoldRefCount=0
+// CHECK: Mapping exists with HstPtrBegin=[[DEVICE_DATA_HST_PTR]], TgtPtrBegin=[[DEVICE_DATA_HST_PTR]], Size=8, DynRefCount=0 (decremented, delayed deletion), HoldRefCount=0
+
+// CHECK: Removing map entry with HstPtrBegin=[[ALLOC_HST_PTR]]{{.*}} Size=4096
+// CHECK: Removing map entry with HstPtrBegin=[[DEVICE_ALLOC_HST_PTR]]{{.*}} Size=8
+// CHECK: Removing map entry with HstPtrBegin=[[DATA_HST_PTR]]{{.*}} Size=4096
+// CHECK: Removing map entry with HstPtrBegin=[[DEVICE_DATA_HST_PTR]]{{.*}} Size=8
+// clang-format on
+
+#pragma omp target map(tofrom : device_data, device_alloc)                     \
+    map(close, to : alloc[ : N], data[ : N])
+  {
+    device_data = &data[0];
+    device_alloc = &alloc[0];
+
+    for (int i = 0; i < N; i++) {
+      alloc[i] += 1;
+      data[i] += 1;
+    }
+  }
+
+  if (device_alloc != host_alloc)
+    printf("Address of alloc on device different from host address.\n");
+
+  if (device_data != host_data)
+    printf("Address of data on device different from host address.\n");
+
+  // On the host, check that the arrays have been updated.
+  fails = 0;
+  for (int i = 0; i < N; i++) {
+    if (alloc[i] != 10)
+      fails++;
+  }
+  printf("Alloc host values not updated: %s\n",
+         (fails == 0) ? "Succeeded" : "Failed");
+
+  fails = 0;
+  for (int i = 0; i < N; i++) {
+    if (data[i] != 1)
+      fails++;
+  }
+  printf("Data host values not updated: %s\n",
+         (fails == 0) ? "Succeeded" : "Failed");
+
+  //
+  // Test that updates on the device are visible on host
+  // when a from is used.
+  //
+
+  for (int i = 0; i < N; i++) {
+    alloc[i] += 1;
+    data[i] += 1;
+  }
+
+// clang-format off
+  // CHECK: Creating new map entry with HstPtrBase=[[ALLOC_HST_PTR:0x.*]], HstPtrBegin=[[ALLOC_HST_PTR]], TgtAllocBegin=[[ALLOC_DEV_PTR:0x.*]], TgtPtrBegin=[[ALLOC_DEV_PTR]], Size=4096, DynRefCount=1, HoldRefCount=0
+  // CHECK: Copying data from host to device, HstPtr=[[ALLOC_HST_PTR]], TgtPtr=[[ALLOC_DEV_PTR]], Size=4096
+
+  // CHECK: Creating new map entry with HstPtrBase=[[DATA_HST_PTR:0x.*]], HstPtrBegin=[[DATA_HST_PTR]], TgtAllocBegin=[[DATA_DEV_PTR:0x.*]], TgtPtrBegin=[[DATA_DEV_PTR]], Size=4096, DynRefCount=1, HoldRefCount=0
+  // CHECK: Copying data from host to device, HstPtr=[[DATA_HST_PTR]], TgtPtr=[[DATA_DEV_PTR]], Size=4096
+
+  // CHECK: Mapping exists with HstPtrBegin=[[ALLOC_HST_PTR]], TgtPtrBegin=[[ALLOC_DEV_PTR]], Size=4096, DynRefCount=1 (update suppressed), HoldRefCount=0
+  // CHECK: Mapping exists with HstPtrBegin=[[DATA_HST_PTR]], TgtPtrBegin=[[DATA_DEV_PTR]], Size=4096, DynRefCount=1 (update suppressed), HoldRefCount=0
+
+  // CHECK: Launching kernel __omp_offloading_{{.*}}_main_l{{.*}} with 1 blocks and 256 threads in Generic mode
+
+  // CHECK: Mapping exists with HstPtrBegin=[[DATA_HST_PTR]], TgtPtrBegin=[[DATA_DEV_PTR]], Size=4096, DynRefCount=0 (decremented, delayed deletion), HoldRefCount=0
+  // CHECK: Mapping exists with HstPtrBegin=[[ALLOC_HST_PTR]], TgtPtrBegin=[[ALLOC_DEV_PTR]], Size=4096, DynRefCount=0 (decremented, delayed deletion), HoldRefCount=0
+
+  // CHECK: Removing map entry with HstPtrBegin=[[DATA_HST_PTR]]{{.*}} Size=4096
+  // CHECK: Removing map entry with HstPtrBegin=[[ALLOC_HST_PTR]]{{.*}} Size=4096
+// clang-format on
+
+  int alloc_fails = 0;
+  int data_fails = 0;
+#pragma omp target map(close, tofrom : alloc[ : N], data[ : N])                \
+    map(tofrom : alloc_fails, data_fails)
+  {
+    for (int i = 0; i < N; i++) {
+      if (alloc[i] != 11)
+        alloc_fails++;
+    }
+    for (int i = 0; i < N; i++) {
+      if (data[i] != 2)
+        data_fails++;
+    }
+
+    // Update values on the device
+    for (int i = 0; i < N; i++) {
+      alloc[i] += 1;
+      data[i] += 1;
+    }
+  }
+
+  printf("Alloc device values are correct: %s\n",
+         (alloc_fails == 0) ? "Succeeded" : "Failed");
+  printf("Data device values are correct: %s\n",
+         (data_fails == 0) ? "Succeeded" : "Failed");
+
+  fails = 0;
+  for (int i = 0; i < N; i++) {
+    if (alloc[i] != 12)
+      fails++;
+  }
+  printf("Alloc host values updated: %s\n",
+         (fails == 0) ? "Succeeded" : "Failed");
+
+  fails = 0;
+  for (int i = 0; i < N; i++) {
+    if (data[i] != 3)
+      fails++;
+  }
+  printf("Data host values updated: %s\n",
+         (fails == 0) ? "Succeeded" : "Failed");
+
+  free(alloc);
+
+  // CHECK: Address of alloc on device different from host address.
+  // CHECK: Address of data on device different from host address.
+  // On the host, check that the arrays have been updated.
+  // CHECK: Alloc host values not updated: Succeeded
+  // CHECK: Data host values not updated: Succeeded
+
+  // CHECK: Alloc device values are correct: Succeeded
+  // CHECK: Data device values are correct: Succeeded
+  // CHECK: Alloc host values updated: Succeeded
+  // CHECK: Data host values updated: Succeeded
+
+  // CHECK: Done!
+  printf("Done!\n");
+
+  return 0;
+}
diff --git a/openmp/libomptarget/test/unified_shared_memory/unified_map_checks_error.cpp b/openmp/libomptarget/test/unified_shared_memory/unified_map_checks_error.cpp
new file mode 100644
index 000000000000000..67ab3702a0d56cf
--- /dev/null
+++ b/openmp/libomptarget/test/unified_shared_memory/unified_map_checks_error.cpp
@@ -0,0 +1,36 @@
+// clang-format off
+// RUN: %libomptarget-compilexx-generic && env HSA_XNACK=1 LIBOMPTARGET_INFO=-1 %libomptarget-run-fail-generic 2>&1 | %fcheck-generic
+// clang-format on
+
+// REQUIRES: amdgcn-amd-amdhsa
+// UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
+
+#include <omp.h>
+
+#include <cassert>
+#include <iostream>
+
+#pragma omp requires unified_shared_memory
+
+int main(int argc, char *argv[]) {
+  int *v = (int *)malloc(sizeof(int) * 10);
+
+// clang-format off
+// CHECK: Entering OpenMP data region with being_mapper at {{.*}} with 1 arguments:
+// CHECK: Creating new map entry ONLY with HstPtrBase=[[V_HST_PTR_ADDR:0x.*]], HstPtrBegin=[[V_HST_PTR_ADDR]], TgtAllocBegin=[[V_HST_PTR_ADDR]], TgtPtrBegin=[[V_HST_PTR_ADDR]], Size=200, DynRefCount=1, HoldRefCount=0
+// CHECK: explicit extension not allowed: host address specified is [[V_HST_PTR_ADDR]] (280 bytes), but device allocation maps to host at [[V_HST_PTR_ADDR]] (200 bytes)
+// CHECK: Call to getTargetPointer returned null pointer (device failure or illegal mapping).
+// clang-format on
+#pragma omp target enter data map(to : v[ : 50])
+
+#pragma omp target enter data map(to : v[ : 70])
+
+#pragma omp target
+  {}
+
+  free(v);
+
+  std::cout << "PASS\n";
+  return 0;
+}
+// CHECK-NOT: PASS
diff --git a/openmp/libomptarget/test/unified_shared_memory/unified_map_checks_no_target.cpp b/openmp/libomptarget/test/unified_shared_memory/unified_map_checks_no_target.cpp
new file mode 100644
index 000000000000000..8983ea994e25c42
--- /dev/null
+++ b/openmp/libomptarget/test/unified_shared_memory/unified_map_checks_no_target.cpp
@@ -0,0 +1,36 @@
+// clang-format off
+// RUN: %libomptarget-compilexx-generic && env HSA_XNACK=1 LIBOMPTARGET_INFO=-1 %libomptarget-run-generic 2>&1 | %fcheck-generic
+// clang-format on
+
+// REQUIRES: amdgcn-amd-amdhsa
+// UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
+
+#include <omp.h>
+
+#include <cassert>
+#include <iostream>
+
+#pragma omp requires unified_shared_memory
+
+/// In the current implementation the lack of a target region in the code
+/// means that unified shared memory is not being enabled even if the pragma
+/// is used explicitly. The code below showcases the copying of data to the
+/// GPU.
+
+int main(int argc, char *argv[]) {
+  int *v = (int *)malloc(sizeof(int) * 10);
+
+// clang-format off
+// CHECK: Entering OpenMP data region with being_mapper at {{.*}} with 1 arguments:
+// CHECK: Creating new map entry with HstPtrBase=[[V_HST_PTR_ADDR:0x.*]], HstPtrBegin=[[V_HST_PTR_ADDR]], TgtAllocBegin=[[V_DEV_PTR_ADDR:0x.*]], TgtPtrBegin=[[V_DEV_PTR_ADDR]], Size=200, DynRefCount=1, HoldRefCount=0
+// CHECK: Copying data from host to device, HstPtr=[[V_HST_PTR_ADDR]], TgtPtr=[[V_DEV_PTR_ADDR]], Size=200
+// clang-format on
+
+#pragma omp target enter data map(to : v[ : 50])
+
+  free(v);
+
+  std::cout << "PASS\n";
+  return 0;
+}
+// CHECK: PASS
diff --git a/openmp/libomptarget/test/unified_shared_memory/unified_map_checks_scalars.cpp b/openmp/libomptarget/test/unified_shared_memory/unified_map_checks_scalars.cpp
new file mode 100644
index 000000000000000..9f7ceacb5791779
--- /dev/null
+++ b/openmp/libomptarget/test/unified_shared_memory/unified_map_checks_scalars.cpp
@@ -0,0 +1,99 @@
+// clang-format off
+// RUN: %libomptarget-compilexx-generic && env HSA_XNACK=1 LIBOMPTARGET_INFO=-1 %libomptarget-run-generic 2>&1 | %fcheck-generic
+// clang-format on
+
+// REQUIRES: amdgcn-amd-amdhsa
+// UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
+
+#include <omp.h>
+
+#include <cassert>
+#include <iostream>
+
+#pragma omp requires unified_shared_memory
+
+int main(int argc, char *argv[]) {
+  int x;
+  int y;
+  int z;
+
+  x = 5;
+  y = 7;
+  z = 11;
+
+  int *v = (int *)malloc(sizeof(int) * 10);
+// clang-format off
+// CHECK: Entering OpenMP data region with being_mapper at {{.*}} with 1 arguments:
+// CHECK: Creating new map entry ONLY with HstPtrBase=[[Z_HST_PTR_BEGIN:0x.*]], HstPtrBegin=[[Z_HST_PTR_BEGIN]], TgtAllocBegin=[[Z_HST_PTR_BEGIN]], TgtPtrBegin=[[Z_HST_PTR_BEGIN]], Size=4, DynRefCount=1, HoldRefCount=0
+// CHECK: OpenMP Host-Device pointer mappings after block
+// CHECK: Host Ptr
+// CHECK: [[Z_HST_PTR_BEGIN]]{{.*}}[[Z_HST_PTR_BEGIN]]
+// clang-format on
+#pragma omp target enter data map(to : z)
+// clang-format off
+// CHECK: Entering OpenMP kernel at {{.*}} with 4 arguments:
+// CHECK: Creating new map entry ONLY with HstPtrBase=[[X_HST_PTR_BEGIN:0x.*]], HstPtrBegin=[[X_HST_PTR_BEGIN]], TgtAllocBegin=[[X_HST_PTR_BEGIN]], TgtPtrBegin=[[X_HST_PTR_BEGIN]], Size=4, DynRefCount=1, HoldRefCount=0
+// CHECK: Creating new map entry ONLY with HstPtrBase=[[Y_HST_PTR_BEGIN:0x.*]], HstPtrBegin=[[Y_HST_PTR_BEGIN]], TgtAllocBegin=[[Y_HST_PTR_BEGIN]], TgtPtrBegin=[[Y_HST_PTR_BEGIN]], Size=4, DynRefCount=1, HoldRefCount=0
+// CHECK: variable {{.*}} does not have a valid device counterpart
+// CHECK: Mapping exists with HstPtrBegin=[[X_HST_PTR_BEGIN]]{{.*}} Size=4, DynRefCount=1 (update suppressed), HoldRefCount=0
+// CHECK: Mapping exists with HstPtrBegin=[[Y_HST_PTR_BEGIN]]{{.*}} Size=4, DynRefCount=1 (update suppressed), HoldRefCount=0
+// CHECK: Launching kernel __omp_offloading_{{.*}}_main_l{{.*}} with 1 blocks and 256 threads in Generic mode
+// clang-format on
+#pragma omp target map(tofrom : x) map(always, tofrom : y) map(to : v[ : 0])
+  {
+    x++;
+    y++;
+    z++;
+  }
+
+// clang-format off
+// CHECK: Mapping exists with HstPtrBegin=[[Y_HST_PTR_BEGIN]]{{.*}} Size=4, DynRefCount=0 (decremented, delayed deletion), HoldRefCount=0
+// CHECK: Mapping exists with HstPtrBegin=[[X_HST_PTR_BEGIN]]{{.*}} Size=4, DynRefCount=0 (decremented, delayed deletion), HoldRefCount=0
+// CHECK: Removing map entry with HstPtrBegin=[[Y_HST_PTR_BEGIN]], TgtPtrBegin=[[Y_HST_PTR_BEGIN]], Size=4
+// CHECK: Removing map entry with HstPtrBegin=[[X_HST_PTR_BEGIN]], TgtPtrBegin=[[X_HST_PTR_BEGIN]], Size=4
+// CHECK: OpenMP Host-Device pointer mappings after block
+// CHECK: Host Ptr
+// CHECK: [[Z_HST_PTR_BEGIN]]{{.*}}[[Z_HST_PTR_BEGIN]]
+// clang-format on
+  printf("x = %d, y = %d, z = %d\n", x, y, z);
+
+// clang-format off
+// CHECK: Entering OpenMP kernel at {{.*}} with 4 arguments:
+// CHECK: Creating new map entry ONLY with HstPtrBase=[[X_HST_PTR_BEGIN:0x.*]], HstPtrBegin=[[X_HST_PTR_BEGIN]], TgtAllocBegin=[[X_HST_PTR_BEGIN]], TgtPtrBegin=[[X_HST_PTR_BEGIN]], Size=4, DynRefCount=1, HoldRefCount=0
+// CHECK: Creating new map entry ONLY with HstPtrBase=[[Y_HST_PTR_BEGIN:0x.*]], HstPtrBegin=[[Y_HST_PTR_BEGIN]], TgtAllocBegin=[[Y_HST_PTR_BEGIN]], TgtPtrBegin=[[Y_HST_PTR_BEGIN]], Size=4, DynRefCount=1, HoldRefCount=0
+// CHECK: variable {{.*}} does not have a valid device counterpart
+// CHECK: Mapping exists with HstPtrBegin=[[X_HST_PTR_BEGIN]]{{.*}} Size=4, DynRefCount=1 (update suppressed), HoldRefCount=0
+// CHECK: Mapping exists with HstPtrBegin=[[Y_HST_PTR_BEGIN]]{{.*}} Size=4, DynRefCount=1 (update suppressed), HoldRefCount=0
+// CHECK: Launching kernel __omp_offloading_{{.*}}_main_l{{.*}} with 1 blocks and 256 threads in Generic mode
+// clang-format on
+#pragma omp target map(tofrom : x) map(always, tofrom : y) map(to : v[ : 0])
+  {
+    x++;
+    y++;
+    z++;
+  }
+// clang-format off
+// CHECK: Mapping exists with HstPtrBegin=[[Y_HST_PTR_BEGIN]]{{.*}} Size=4, DynRefCount=0 (decremented, delayed deletion), HoldRefCount=0
+// CHECK: Mapping exists with HstPtrBegin=[[X_HST_PTR_BEGIN]]{{.*}} Size=4, DynRefCount=0 (decremented, delayed deletion), HoldRefCount=0
+// CHECK: Removing map entry with HstPtrBegin=[[Y_HST_PTR_BEGIN]], TgtPtrBegin=[[Y_HST_PTR_BEGIN]], Size=4
+// CHECK: Removing map entry with HstPtrBegin=[[X_HST_PTR_BEGIN]], TgtPtrBegin=[[X_HST_PTR_BEGIN]], Size=4
+// CHECK: OpenMP Host-Device pointer mappings after block
+// CHECK: Host Ptr
+// CHECK: [[Z_HST_PTR_BEGIN]]{{.*}}[[Z_HST_PTR_BEGIN]]
+// clang-format on
+#pragma omp target exit data map(from : z)
+// clang-format off
+// CHECK: Exiting OpenMP data region with end_mapper at {{.*}} with 1 arguments:
+// CHECK: Mapping exists with HstPtrBegin=[[Z_HST_PTR_BEGIN]], TgtPtrBegin=[[Z_HST_PTR_BEGIN]], Size=4, DynRefCount=0 (decremented, delayed deletion), HoldRefCount=0
+// CHECK: Removing map entry with HstPtrBegin=[[Z_HST_PTR_BEGIN]], TgtPtrBegin=[[Z_HST_PTR_BEGIN]], Size=4
+// clang-format on
+  printf("x = %d, y = %d, z = %d\n", x, y, z);
+
+  free(v);
+
+  std::cout << "PASS\n";
+  return 0;
+}
+// CHECK: x = 6, y = 8, z = 11
+// CHECK: x = 7, y = 9, z = 11
+// CHECK: PASS
diff --git a/openmp/libomptarget/test/unified_shared_memory/unified_map_checks_shared_update.cpp b/openmp/libomptarget/test/unified_shared_memory/unified_map_checks_shared_update.cpp
new file mode 100644
index 000000000000000..01aa084afb8b17a
--- /dev/null
+++ b/openmp/libomptarget/test/unified_shared_memory/unified_map_checks_shared_update.cpp
@@ -0,0 +1,137 @@
+// clang-format off
+// RUN: %libomptarget-compilexx-generic && env HSA_XNACK=1 LIBOMPTARGET_INFO=-1 %libomptarget-run-generic 2>&1 | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
+
+// REQUIRES: amdgcn-amd-amdhsa
+
+#include <omp.h>
+#include <stdio.h>
+
+#pragma omp requires unified_shared_memory
+
+#define N 1024
+
+int main(int argc, char *argv[]) {
+  int fails;
+  void *host_alloc, *device_alloc;
+  void *host_data, *device_data;
+  int *alloc = (int *)malloc(N * sizeof(int));
+  int data[N];
+
+  for (int i = 0; i < N; ++i) {
+    alloc[i] = 10;
+    data[i] = 1;
+  }
+
+  host_data = &data[0];
+  host_alloc = &alloc[0];
+
+// clang-format off
+// CHECK: Creating new map entry ONLY with HstPtrBase=[[DEVICE_DATA_HST_PTR:0x.*]], HstPtrBegin=[[DEVICE_DATA_HST_PTR]], TgtAllocBegin=[[DEVICE_DATA_HST_PTR]], TgtPtrBegin=[[DEVICE_DATA_HST_PTR]], Size=8, DynRefCount=1, HoldRefCount=0
+// CHECK: Creating new map entry ONLY with HstPtrBase=[[DATA_HST_PTR:0x.*]], HstPtrBegin=[[DATA_HST_PTR]], TgtAllocBegin=[[DATA_HST_PTR]], TgtPtrBegin=[[DATA_HST_PTR]], Size=4096, DynRefCount=1, HoldRefCount=0
+// CHECK: Creating new map entry ONLY with HstPtrBase=[[DEVICE_ALLOC_HST_PTR:0x.*]], HstPtrBegin=[[DEVICE_ALLOC_HST_PTR]], TgtAllocBegin=[[DEVICE_ALLOC_HST_PTR]], TgtPtrBegin=[[DEVICE_ALLOC_HST_PTR]], Size=8, DynRefCount=1, HoldRefCount=0
+// CHECK: Mapping exists with HstPtrBegin=[[DEVICE_DATA_HST_PTR]], TgtPtrBegin=[[DEVICE_DATA_HST_PTR]], Size=8, DynRefCount=1 (update suppressed), HoldRefCount=0
+// CHECK: Mapping exists with HstPtrBegin=[[DATA_HST_PTR]], TgtPtrBegin=[[DATA_HST_PTR]], Size=4096, DynRefCount=1 (update suppressed), HoldRefCount=0
+// CHECK: Mapping exists with HstPtrBegin=[[DEVICE_ALLOC_HST_PTR]], TgtPtrBegin=[[DEVICE_ALLOC_HST_PTR]], Size=8, DynRefCount=1 (update suppressed), HoldRefCount=0
+
+// CHECK: Launching kernel __omp_offloading_{{.*}}_main_l{{.*}} with 1 blocks and 256 threads in Generic mode
+
+// CHECK: Mapping exists with HstPtrBegin=[[DEVICE_ALLOC_HST_PTR]], TgtPtrBegin=[[DEVICE_ALLOC_HST_PTR]], Size=8, DynRefCount=0 (decremented, delayed deletion), HoldRefCount=0
+// CHECK: Mapping exists with HstPtrBegin=[[DATA_HST_PTR]], TgtPtrBegin=[[DATA_HST_PTR]], Size=4096, DynRefCount=0 (decremented, delayed deletion), HoldRefCount=0
+// CHECK: Mapping exists with HstPtrBegin=[[DEVICE_DATA_HST_PTR]], TgtPtrBegin=[[DEVICE_DATA_HST_PTR]], Size=8, DynRefCount=0 (decremented, delayed deletion), HoldRefCount=0
+
+// CHECK: Removing map entry with HstPtrBegin=[[DEVICE_ALLOC_HST_PTR]]{{.*}} Size=8
+// CHECK: Removing map entry with HstPtrBegin=[[DATA_HST_PTR]]{{.*}} Size=4096
+// CHECK: Removing map entry with HstPtrBegin=[[DEVICE_DATA_HST_PTR]]{{.*}} Size=8
+// clang-format on
+
+// implicit mapping of data
+#pragma omp target map(tofrom : device_data, device_alloc)
+  {
+    device_data = &data[0];
+    device_alloc = &alloc[0];
+
+    for (int i = 0; i < N; i++) {
+      alloc[i] += 1;
+      data[i] += 1;
+    }
+  }
+
+  if (device_alloc == host_alloc)
+    printf("Address of alloc on device matches host address.\n");
+
+  if (device_data == host_data)
+    printf("Address of data on device matches host address.\n");
+
+  // On the host, check that the arrays have been updated.
+  fails = 0;
+  for (int i = 0; i < N; i++) {
+    if (alloc[i] != 11)
+      fails++;
+  }
+  printf("Alloc device values updated: %s\n",
+         (fails == 0) ? "Succeeded" : "Failed");
+
+  fails = 0;
+  for (int i = 0; i < N; i++) {
+    if (data[i] != 2)
+      fails++;
+  }
+  printf("Data device values updated: %s\n",
+         (fails == 0) ? "Succeeded" : "Failed");
+
+  //
+  // Test that updates on the host and on the device are both visible.
+  //
+
+  // Update on the host.
+  for (int i = 0; i < N; ++i) {
+    alloc[i] += 1;
+    data[i] += 1;
+  }
+
+// clang-format off
+// CHECK: Creating new map entry ONLY with HstPtrBase=[[DATA_HST_PTR]], HstPtrBegin=[[DATA_HST_PTR]], TgtAllocBegin=[[DATA_HST_PTR]], TgtPtrBegin=[[DATA_HST_PTR]], Size=4096, DynRefCount=1, HoldRefCount=0
+
+// CHECK: Mapping exists with HstPtrBegin=[[DATA_HST_PTR]], TgtPtrBegin=[[DATA_HST_PTR]], Size=4096, DynRefCount=1 (update suppressed), HoldRefCount=0
+
+// CHECK: Launching kernel __omp_offloading_{{.*}}_main_l{{.*}} with 1 blocks and 256 threads in Generic mode
+
+// CHECK: Mapping exists with HstPtrBegin=[[DATA_HST_PTR]], TgtPtrBegin=[[DATA_HST_PTR]], Size=4096, DynRefCount=0 (decremented, delayed deletion), HoldRefCount=0
+
+// CHECK: Removing map entry with HstPtrBegin=[[DATA_HST_PTR]]{{.*}} Size=4096
+// clang-format on
+
+  int alloc_fails = 0;
+  int data_fails = 0;
+#pragma omp target
+  {
+    for (int i = 0; i < N; i++) {
+      if (alloc[i] != 12)
+        alloc_fails++;
+    }
+    for (int i = 0; i < N; i++) {
+      if (data[i] != 3)
+        data_fails++;
+    }
+  }
+  printf("Alloc host values updated: %s\n",
+         (alloc_fails == 0) ? "Succeeded" : "Failed");
+  printf("Data host values updated: %s\n",
+         (data_fails == 0) ? "Succeeded" : "Failed");
+  free(alloc);
+
+  // CHECK: Address of alloc on device matches host address.
+  // CHECK: Address of data on device matches host address.
+  // CHECK: Alloc device values updated: Succeeded
+  // CHECK: Data device values updated: Succeeded
+
+  // CHECK: Alloc host values updated: Succeeded
+  // CHECK: Data host values updated: Succeeded
+
+  printf("Done!\n");
+
+  return 0;
+}
diff --git a/openmp/libomptarget/test/unified_shared_memory/zero_sized_array.cpp b/openmp/libomptarget/test/unified_shared_memory/zero_sized_array.cpp
new file mode 100644
index 000000000000000..e24a653e61fe74c
--- /dev/null
+++ b/openmp/libomptarget/test/unified_shared_memory/zero_sized_array.cpp
@@ -0,0 +1,31 @@
+// clang-format off
+// RUN: %libomptarget-compilexx-generic && env HSA_XNACK=1 LIBOMPTARGET_INFO=-1 %libomptarget-run-generic 2>&1 | %fcheck-generic
+// clang-format on
+
+// REQUIRES: amdgcn-amd-amdhsa
+// UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
+
+#include <omp.h>
+
+#include <cassert>
+#include <iostream>
+
+#pragma omp requires unified_shared_memory
+
+int main(int argc, char *argv[]) {
+  int *v = (int *)malloc(sizeof(int) * 10);
+
+  printf("host address of v = %p\n", v);
+
+// CHECK: variable {{.*}} does not have a valid device counterpart
+#pragma omp target map(to : v[ : 0])
+  { printf("device address of v = %p\n", v); }
+
+  free(v);
+
+  std::cout << "PASS\n";
+  return 0;
+}
+// CHECK: host address of v = [[ADDR_OF_V:0x.*]]
+// TODO: once printf is supported add check for ADDR_OF_V on device
+// CHECK: PASS



More information about the flang-commits mailing list