[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