[clang] [llvm] [WIP][OFFLOAD][OpenMP] Fix USM close mapping, and track declare-target storage under USM (PR #213214)
Abhinav Gaba via llvm-commits
llvm-commits at lists.llvm.org
Tue Aug 4 22:26:59 PDT 2026
https://github.com/abhinavgaba updated https://github.com/llvm/llvm-project/pull/213214
>From 9add159bdad5f6783443df5bdf135b9a1c11b3ba Mon Sep 17 00:00:00 2001
From: "Fine, Gregory" <gregory.fine at intel.com>
Date: Tue, 7 Jul 2026 16:52:13 -0700
Subject: [PATCH 1/9] Fix an issue with nested close, alloc mapping when using
USM
---
offload/libomptarget/OpenMP/Mapping.cpp | 8 +++++
.../unified_shared_memory/close_sub_section.c | 32 +++++++++++++++++++
.../close_sub_section_oob.c | 32 +++++++++++++++++++
3 files changed, 72 insertions(+)
create mode 100644 offload/test/unified_shared_memory/close_sub_section.c
create mode 100644 offload/test/unified_shared_memory/close_sub_section_oob.c
diff --git a/offload/libomptarget/OpenMP/Mapping.cpp b/offload/libomptarget/OpenMP/Mapping.cpp
index 1bb2e424bd083..a9266bf7ca1f9 100644
--- a/offload/libomptarget/OpenMP/Mapping.cpp
+++ b/offload/libomptarget/OpenMP/Mapping.cpp
@@ -282,6 +282,14 @@ TargetPointerResultTy MappingInfoTy::getTargetPointer(
LR.TPR.Flags.IsPresent = false;
LR.TPR.Flags.IsHostPointer = true;
LR.TPR.TargetPointer = HstPtrBegin;
+ // Create a mapping for a case when map(close, alloc:...) is applied to a
+ // subsection of previously mapped allocation. The mapping would prevent
+ // map(close, alloc:...) from creating a new allocation as it would reuse
+ // the mapped allocation instead.
+ HDTTMap->emplace(new HostDataToTargetTy(
+ (uintptr_t)HstPtrBase, (uintptr_t)HstPtrBegin,
+ (uintptr_t)HstPtrBegin + Size, (uintptr_t)HstPtrBegin,
+ (uintptr_t)HstPtrBegin, true, HstPtrName));
}
} else if (HasPresentModifier) {
ODBG(ODT_Mapping) << "Mapping required by 'present' map type modifier does "
diff --git a/offload/test/unified_shared_memory/close_sub_section.c b/offload/test/unified_shared_memory/close_sub_section.c
new file mode 100644
index 0000000000000..06cf444c49283
--- /dev/null
+++ b/offload/test/unified_shared_memory/close_sub_section.c
@@ -0,0 +1,32 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+// REQUIRES: unified_shared_memory
+// UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
+
+// amdgpu runtime crash
+// Fails on nvptx with error: an illegal memory access was encountered
+// UNSUPPORTED: amdgcn-amd-amdhsa
+// UNSUPPORTED: nvptx64-nvidia-cuda
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: intelgpu
+
+#pragma omp requires unified_shared_memory
+
+#include <stdio.h>
+
+int main() {
+ double base[10] = {0};
+
+#pragma omp target data map(from : base[0 : 10])
+ {
+// close range covers base[1] and should properly
+// update the budder from the outer mapping.
+#pragma omp target map(close, alloc : base[1 : 9])
+ {
+ base[1] = 99.0;
+ }
+ }
+ // CHECK: base[1] = 99.000000 (expected 99.0)
+ printf("base[1] = %f (expected 99.0)\n", base[1]);
+ return 0;
+}
\ No newline at end of file
diff --git a/offload/test/unified_shared_memory/close_sub_section_oob.c b/offload/test/unified_shared_memory/close_sub_section_oob.c
new file mode 100644
index 0000000000000..1638f4a6f3c4c
--- /dev/null
+++ b/offload/test/unified_shared_memory/close_sub_section_oob.c
@@ -0,0 +1,32 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+// REQUIRES: unified_shared_memory
+// UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
+
+// amdgpu runtime crash
+// Fails on nvptx with error: an illegal memory access was encountered
+// UNSUPPORTED: amdgcn-amd-amdhsa
+// UNSUPPORTED: nvptx64-nvidia-cuda
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: intelgpu
+
+#pragma omp requires unified_shared_memory
+
+#include <stdio.h>
+
+int main() {
+ double base[10] = {0};
+
+#pragma omp target data map(from : base[0 : 10])
+ {
+// close on sub-section base[2:8]; access to base[1] is outside the close
+// buffer
+#pragma omp target map(close, alloc : base[2 : 8])
+ {
+ base[1] = 99.0;
+ }
+ }
+ // CHECK: base[1] = 99.000000 (expected 99.0)
+ printf("base[1] = %f (expected 99.0)\n", base[1]);
+ return 0;
+}
\ No newline at end of file
>From 646b2dd87aed73f39b1c0e0382e1260ce8a1451c Mon Sep 17 00:00:00 2001
From: "Fine, Gregory" <gregory.fine at intel.com>
Date: Wed, 8 Jul 2026 15:33:38 -0700
Subject: [PATCH 2/9] Add test
---
.../unified_shared_memory/check_tracking.c | 30 +++++++++++++++++++
1 file changed, 30 insertions(+)
create mode 100644 offload/test/unified_shared_memory/check_tracking.c
diff --git a/offload/test/unified_shared_memory/check_tracking.c b/offload/test/unified_shared_memory/check_tracking.c
new file mode 100644
index 0000000000000..7607e9df90f87
--- /dev/null
+++ b/offload/test/unified_shared_memory/check_tracking.c
@@ -0,0 +1,30 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+// REQUIRES: unified_shared_memory
+// UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
+
+// amdgpu runtime crash
+// Fails on nvptx with error: an illegal memory access was encountered
+// UNSUPPORTED: amdgcn-amd-amdhsa
+// UNSUPPORTED: nvptx64-nvidia-cuda
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: intelgpu
+
+#pragma omp requires unified_shared_memory
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+ int x = 111;
+
+ // CHECK: present when unmapped: 0
+ printf("present when unmapped: %d\n",
+ omp_target_is_present(&x, omp_get_default_device()));
+
+#pragma omp target_enter_data map(alloc : x)
+
+ // CHECK: present after mapping: 1
+ printf("present after mapping: %d\n",
+ omp_target_is_present(&x, omp_get_default_device()));
+ return 0;
+}
\ No newline at end of file
>From 15bbe4e9030dc2f35a7ee210efc917031a55d205 Mon Sep 17 00:00:00 2001
From: "Fine, Gregory" <gregory.fine at intel.com>
Date: Thu, 9 Jul 2026 16:18:41 -0700
Subject: [PATCH 3/9] Add handling of global base
---
offload/libomptarget/OpenMP/Mapping.cpp | 2 +-
.../close_sub_section_global.c | 33 +++++++++++++++++++
.../close_sub_section_global_oob.c | 33 +++++++++++++++++++
3 files changed, 67 insertions(+), 1 deletion(-)
create mode 100644 offload/test/unified_shared_memory/close_sub_section_global.c
create mode 100644 offload/test/unified_shared_memory/close_sub_section_global_oob.c
diff --git a/offload/libomptarget/OpenMP/Mapping.cpp b/offload/libomptarget/OpenMP/Mapping.cpp
index a9266bf7ca1f9..064df0a0c7722 100644
--- a/offload/libomptarget/OpenMP/Mapping.cpp
+++ b/offload/libomptarget/OpenMP/Mapping.cpp
@@ -263,7 +263,7 @@ TargetPointerResultTy MappingInfoTy::getTargetPointer(
"exist for host address " DPxMOD " (%" PRId64 " bytes)",
DPxPTR(HstPtrBegin), Size);
} else if ((PM->getRequirements() & OMP_REQ_UNIFIED_SHARED_MEMORY &&
- !HasCloseModifier) ||
+ (!HasCloseModifier || LR.TPR.getEntry() != nullptr)) ||
(PM->getRequirements() & OMPX_REQ_AUTO_ZERO_COPY)) {
// If unified shared memory is active, implicitly mapped variables that are
diff --git a/offload/test/unified_shared_memory/close_sub_section_global.c b/offload/test/unified_shared_memory/close_sub_section_global.c
new file mode 100644
index 0000000000000..7bf128d636fa8
--- /dev/null
+++ b/offload/test/unified_shared_memory/close_sub_section_global.c
@@ -0,0 +1,33 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+// REQUIRES: unified_shared_memory
+// UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
+
+// amdgpu runtime crash
+// Fails on nvptx with error: an illegal memory access was encountered
+// UNSUPPORTED: amdgcn-amd-amdhsa
+// UNSUPPORTED: nvptx64-nvidia-cuda
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: intelgpu
+
+#pragma omp requires unified_shared_memory
+
+#include <stdio.h>
+
+#pragma omp begin declare target
+double base[10] = {0};
+#pragma omp end declare target
+
+int main() {
+
+// close range covers base[1] and should properly
+// update the budder from the outer mapping.
+#pragma omp target map(close, alloc : base[1 : 9])
+ {
+ base[1] = 99.0;
+ }
+#pragma omp target update from(base)
+ // CHECK: base[1] = 99.000000 (expected 99.0)
+ printf("base[1] = %f (expected 99.0)\n", base[1]);
+ return 0;
+}
\ No newline at end of file
diff --git a/offload/test/unified_shared_memory/close_sub_section_global_oob.c b/offload/test/unified_shared_memory/close_sub_section_global_oob.c
new file mode 100644
index 0000000000000..eaf5ff8c74490
--- /dev/null
+++ b/offload/test/unified_shared_memory/close_sub_section_global_oob.c
@@ -0,0 +1,33 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+// REQUIRES: unified_shared_memory
+// UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
+
+// amdgpu runtime crash
+// Fails on nvptx with error: an illegal memory access was encountered
+// UNSUPPORTED: amdgcn-amd-amdhsa
+// UNSUPPORTED: nvptx64-nvidia-cuda
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: intelgpu
+
+#pragma omp requires unified_shared_memory
+
+#include <stdio.h>
+
+#pragma omp begin declare target
+double base[10] = {0};
+#pragma omp end declare target
+
+int main() {
+
+// close on sub-section base[2:8]; access to base[1] is outside the close
+// buffer
+#pragma omp target map(close, alloc : base[2 : 8])
+ {
+ base[1] = 99.0;
+ }
+#pragma omp target update from(base)
+ // CHECK: base[1] = 99.000000 (expected 99.0)
+ printf("base[1] = %f (expected 99.0)\n", base[1]);
+ return 0;
+}
\ No newline at end of file
>From 0c5ce50c77cd2966296e451a2c652e33e1c99ab0 Mon Sep 17 00:00:00 2001
From: "Fine, Gregory" <gregory.fine at intel.com>
Date: Mon, 27 Jul 2026 20:38:40 -0700
Subject: [PATCH 4/9] Fix reference counting
---
offload/libomptarget/OpenMP/Mapping.cpp | 34 +++++++++++++++---
offload/libomptarget/PluginManager.cpp | 22 ++++++++++++
.../unified_shared_memory/check_tracking.c | 35 +++++++++++++++----
3 files changed, 80 insertions(+), 11 deletions(-)
diff --git a/offload/libomptarget/OpenMP/Mapping.cpp b/offload/libomptarget/OpenMP/Mapping.cpp
index 064df0a0c7722..b2b04dd8716b9 100644
--- a/offload/libomptarget/OpenMP/Mapping.cpp
+++ b/offload/libomptarget/OpenMP/Mapping.cpp
@@ -263,7 +263,18 @@ TargetPointerResultTy MappingInfoTy::getTargetPointer(
"exist for host address " DPxMOD " (%" PRId64 " bytes)",
DPxPTR(HstPtrBegin), Size);
} else if ((PM->getRequirements() & OMP_REQ_UNIFIED_SHARED_MEMORY &&
- (!HasCloseModifier || LR.TPR.getEntry() != nullptr)) ||
+ (!HasCloseModifier ||
+ // A close mapping should not incur a new
+ // allocation under USM when it is already
+ // "present" on the device. That can either
+ // be due to an overlap with a previously
+ // encountered map (with/without the close
+ // modifier), or it being declare_target
+ // (infinite ref-count).
+ (LR.TPR.getEntry() != nullptr &&
+ (LR.Flags.IsContained || LR.Flags.ExtendsBefore ||
+ LR.Flags.ExtendsAfter ||
+ LR.TPR.getEntry()->isDynRefCountInf())))) ||
(PM->getRequirements() & OMPX_REQ_AUTO_ZERO_COPY)) {
// If unified shared memory is active, implicitly mapped variables that are
@@ -286,10 +297,15 @@ TargetPointerResultTy MappingInfoTy::getTargetPointer(
// subsection of previously mapped allocation. The mapping would prevent
// map(close, alloc:...) from creating a new allocation as it would reuse
// the mapped allocation instead.
- HDTTMap->emplace(new HostDataToTargetTy(
- (uintptr_t)HstPtrBase, (uintptr_t)HstPtrBegin,
- (uintptr_t)HstPtrBegin + Size, (uintptr_t)HstPtrBegin,
- (uintptr_t)HstPtrBegin, true, HstPtrName));
+ LR.TPR.setEntry(
+ HDTTMap
+ ->emplace(new HostDataToTargetTy(
+ (uintptr_t)HstPtrBase, (uintptr_t)HstPtrBegin,
+ (uintptr_t)HstPtrBegin + Size, (uintptr_t)HstPtrBegin,
+ (uintptr_t)HstPtrBegin, HasHoldModifier, HstPtrName))
+ .first->HDTT);
+ if (Device.notifyDataMapped(HstPtrBegin, Size))
+ return TargetPointerResultTy{};
}
} else if (HasPresentModifier) {
ODBG(ODT_Mapping) << "Mapping required by 'present' map type modifier does "
@@ -446,6 +462,14 @@ TargetPointerResultTy MappingInfoTy::getTgtPtrBegin(
"expected correct IsLast prediction for reset");
}
+ if (PM->getRequirements() & OMP_REQ_UNIFIED_SHARED_MEMORY ||
+ PM->getRequirements() & OMPX_REQ_AUTO_ZERO_COPY) {
+ LR.TPR.Flags.IsHostPointer = true;
+ if (!LR.TPR.getEntry()->getTotalRefCount()) {
+ LR.TPR.Flags.IsPresent = false;
+ }
+ }
+
// Increment the number of threads that is using the entry on a
// targetDataEnd, tracking the number of possible "deleters". A thread may
// come to own the entry deletion even if it was not the last one querying
diff --git a/offload/libomptarget/PluginManager.cpp b/offload/libomptarget/PluginManager.cpp
index 41b653a60adfd..d6a27537aedf9 100644
--- a/offload/libomptarget/PluginManager.cpp
+++ b/offload/libomptarget/PluginManager.cpp
@@ -516,6 +516,28 @@ static int loadImagesOntoDevice(DeviceTy &Device) {
CurrDeviceEntryAddr = DevPtr;
}
+ if (PM->getRequirements() & OMP_REQ_UNIFIED_SHARED_MEMORY ||
+ PM->getRequirements() & OMPX_REQ_AUTO_ZERO_COPY) {
+ AsyncInfoTy AsyncInfo(Device);
+ void *DevPtr;
+ Device.retrieveData(&DevPtr, CurrDeviceEntryAddr, sizeof(void *),
+ AsyncInfo, /*Entry=*/nullptr, &HDTTMap);
+ if (AsyncInfo.synchronize() != OFFLOAD_SUCCESS)
+ return OFFLOAD_FAIL;
+
+ ODBG(ODT_Mapping)
+ << "Add mapping from host " << DevPtr << " to device " << DevPtr
+ << " with size " << CurrDeviceEntry->Size;
+
+ HDTTMap->emplace(new HostDataToTargetTy(
+ (uintptr_t)DevPtr /*HstPtrBase*/,
+ (uintptr_t)DevPtr /*HstPtrBegin*/,
+ (uintptr_t)DevPtr + CurrHostEntry->Size /*HstPtrEnd*/,
+ (uintptr_t)DevPtr /*TgtAllocBegin*/,
+ (uintptr_t)DevPtr /*TgtPtrBegin*/, false /*UseHoldRefCount*/,
+ nullptr, true /*IsRefCountINF*/));
+ }
+
ODBG(ODT_Mapping) << "Add mapping from host " << CurrHostEntry->Address
<< " to device " << CurrDeviceEntry->Address
<< " with size " << CurrDeviceEntry->Size
diff --git a/offload/test/unified_shared_memory/check_tracking.c b/offload/test/unified_shared_memory/check_tracking.c
index 7607e9df90f87..808c6befac52b 100644
--- a/offload/test/unified_shared_memory/check_tracking.c
+++ b/offload/test/unified_shared_memory/check_tracking.c
@@ -14,17 +14,40 @@
#include <omp.h>
#include <stdio.h>
+#pragma omp begin declare target
+int x = 111;
+#pragma omp end declare target
+int y = 111;
+
+int present(void *p) {
+ return omp_target_is_present(p, omp_get_default_device());
+}
+
int main() {
- int x = 111;
+ int xl = 111;
// CHECK: present when unmapped: 0
- printf("present when unmapped: %d\n",
- omp_target_is_present(&x, omp_get_default_device()));
+ printf("present when unmapped: %d\n", present(&xl));
-#pragma omp target_enter_data map(alloc : x)
+#pragma omp target_enter_data map(alloc : xl)
// CHECK: present after mapping: 1
- printf("present after mapping: %d\n",
- omp_target_is_present(&x, omp_get_default_device()));
+ printf("present after mapping: %d\n", present(&xl));
+#pragma omp target_exit_data map(from : xl)
+ // CHECK: present after mapping: 0
+ printf("present after mapping: %d\n", present(&xl));
+
+ // CHECK: present when unmapped: 1 0
+ printf("present when unmapped: %d %d\n", present(&x), present(&y));
+
+#pragma omp target_enter_data map(to : x, y)
+
+ // CHECK: present after mapping: 1 1
+ printf("present after mapping: %d %d\n", present(&x), present(&y));
+
+#pragma omp target_exit_data map(from : x, y)
+
+ // CHECK: present after mapping: 1 0
+ printf("present after mapping: %d %d\n", present(&x), present(&y));
return 0;
}
\ No newline at end of file
>From aab9395010b8ac75a92f6be7af6c635fb146c3aa Mon Sep 17 00:00:00 2001
From: "Fine, Gregory" <gregory.fine at intel.com>
Date: Mon, 27 Jul 2026 20:53:36 -0700
Subject: [PATCH 5/9] Add test
---
offload/test/unified_shared_memory/close.c | 35 ++++++++++++++++++++++
1 file changed, 35 insertions(+)
create mode 100644 offload/test/unified_shared_memory/close.c
diff --git a/offload/test/unified_shared_memory/close.c b/offload/test/unified_shared_memory/close.c
new file mode 100644
index 0000000000000..c813663bfbcec
--- /dev/null
+++ b/offload/test/unified_shared_memory/close.c
@@ -0,0 +1,35 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+// REQUIRES: unified_shared_memory
+// UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
+
+// amdgpu runtime crash
+// Fails on nvptx with error: an illegal memory access was encountered
+// UNSUPPORTED: amdgcn-amd-amdhsa
+// UNSUPPORTED: nvptx64-nvidia-cuda
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: intelgpu
+
+// This is an implementation-specific test (it does not reflect official OpenMP
+// expectations). The test is to ensure that we do the optimization to
+// "allocate" a new device copy for maps with `close` modifier, but only when a
+// matching entry doesn't already exist on the device.
+
+#include <assert.h>
+
+int x, y;
+#pragma omp requires unified_shared_memory
+
+int main() {
+ int *xaddr_device, *yaddr_device;
+#pragma omp target data map(tofrom : x)
+#pragma omp target map(close, tofrom : x, y) \
+ map(from : xaddr_device, yaddr_device)
+ {
+ xaddr_device = &x;
+ yaddr_device = &y;
+ }
+
+ assert(xaddr_device == &x && "Mapped variable should not allocate on close.");
+ assert(yaddr_device != &y && "Unmapped variable should allocate on close.");
+}
>From 7a383c7cbc74e9b11f0500f9c201ca78bf7ac001 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Thu, 30 Jul 2026 23:46:46 -0700
Subject: [PATCH 6/9] [OFFLOAD] Fix USM mapping tracking without disabling
device data motion
Reporting every mapping found under unified shared memory as a host pointer also
disables the work a close mapping depends on: the FROM copy-back, target update
and pointer attachment all skip their work for a host pointer. Device writes were
silently dropped, regressing close_modifier.c, close_member.c and associate_ptr.c.
Instead let the reuse entry recorded on the USM host path be reclaimed with the
region that created it, and skip only the device deallocation for it, since its
allocation maps to the host address. That also fixes close_enter_exit.c, where the
entry used to linger with a zero reference count and make a later
map(close, ...) of the same storage look already present.
Gate the close reuse on the containment flags alone. An infinite reference count
is not a usable signal: it is evaluated on whatever entry the lookup landed on, so
it depended on address layout and an unrelated sibling map could displace it.
Adds close_data_motion.c, covering FROM copy-back and target update for a close
mapping under USM, and a CHECK line to close.c, whose RUN line pipes into
FileCheck but which had none.
The declare-target parts are left out. Under USM the offload entry for such a
variable describes the device reference pointer, so the runtime has neither the
variable's address nor its extent and cannot register its storage correctly.
That needs a code-generation change, and the tests that depend on it go with it.
Co-Authored-By: Claude Opus 5 (1M context) <noreply at anthropic.com>
---
offload/libomptarget/OpenMP/Mapping.cpp | 37 ++++-----
offload/libomptarget/PluginManager.cpp | 22 ------
offload/libomptarget/omptarget.cpp | 10 ++-
.../unified_shared_memory/check_tracking.c | 35 ++------
offload/test/unified_shared_memory/close.c | 5 ++
.../unified_shared_memory/close_data_motion.c | 79 +++++++++++++++++++
.../close_sub_section_global.c | 33 --------
.../close_sub_section_global_oob.c | 33 --------
8 files changed, 118 insertions(+), 136 deletions(-)
create mode 100644 offload/test/unified_shared_memory/close_data_motion.c
delete mode 100644 offload/test/unified_shared_memory/close_sub_section_global.c
delete mode 100644 offload/test/unified_shared_memory/close_sub_section_global_oob.c
diff --git a/offload/libomptarget/OpenMP/Mapping.cpp b/offload/libomptarget/OpenMP/Mapping.cpp
index b2b04dd8716b9..295447387c607 100644
--- a/offload/libomptarget/OpenMP/Mapping.cpp
+++ b/offload/libomptarget/OpenMP/Mapping.cpp
@@ -264,17 +264,21 @@ TargetPointerResultTy MappingInfoTy::getTargetPointer(
DPxPTR(HstPtrBegin), Size);
} else if ((PM->getRequirements() & OMP_REQ_UNIFIED_SHARED_MEMORY &&
(!HasCloseModifier ||
- // A close mapping should not incur a new
- // allocation under USM when it is already
- // "present" on the device. That can either
- // be due to an overlap with a previously
- // encountered map (with/without the close
- // modifier), or it being declare_target
- // (infinite ref-count).
+ // A close mapping should not incur a new allocation under USM
+ // when the storage it refers to is already on the device, i.e.
+ // when it overlaps a previously encountered map, with or without
+ // the close modifier.
+ //
+ // Storage that is present for the whole program, such as a
+ // declare-target variable, should be covered by this as well but
+ // is not: under USM the entry registered for such a variable
+ // describes the device reference pointer rather than the
+ // variable, so a mapping of the variable does not overlap it.
+ // Communicating the variable's extent is a code-generation
+ // change.
(LR.TPR.getEntry() != nullptr &&
(LR.Flags.IsContained || LR.Flags.ExtendsBefore ||
- LR.Flags.ExtendsAfter ||
- LR.TPR.getEntry()->isDynRefCountInf())))) ||
+ LR.Flags.ExtendsAfter)))) ||
(PM->getRequirements() & OMPX_REQ_AUTO_ZERO_COPY)) {
// If unified shared memory is active, implicitly mapped variables that are
@@ -462,14 +466,6 @@ TargetPointerResultTy MappingInfoTy::getTgtPtrBegin(
"expected correct IsLast prediction for reset");
}
- if (PM->getRequirements() & OMP_REQ_UNIFIED_SHARED_MEMORY ||
- PM->getRequirements() & OMPX_REQ_AUTO_ZERO_COPY) {
- LR.TPR.Flags.IsHostPointer = true;
- if (!LR.TPR.getEntry()->getTotalRefCount()) {
- LR.TPR.Flags.IsPresent = false;
- }
- }
-
// Increment the number of threads that is using the entry on a
// targetDataEnd, tracking the number of possible "deleters". A thread may
// come to own the entry deletion even if it was not the last one querying
@@ -577,7 +573,12 @@ int MappingInfoTy::deallocTgtPtrAndEntry(HostDataToTargetTy *Entry,
return OFFLOAD_FAIL;
}
- int Ret = Device.deleteData((void *)Entry->TgtAllocBegin);
+ // The reuse entry recorded on the unified-shared-memory host path owns no
+ // device allocation: its allocation maps to the host address itself, so it
+ // must not be handed to deleteData().
+ int Ret = OFFLOAD_SUCCESS;
+ if (Entry->TgtAllocBegin != Entry->HstPtrBegin)
+ Ret = Device.deleteData((void *)Entry->TgtAllocBegin);
// Notify the plugin about the unmapped memory.
Ret |= Device.notifyDataUnmapped((void *)Entry->HstPtrBegin);
diff --git a/offload/libomptarget/PluginManager.cpp b/offload/libomptarget/PluginManager.cpp
index d6a27537aedf9..41b653a60adfd 100644
--- a/offload/libomptarget/PluginManager.cpp
+++ b/offload/libomptarget/PluginManager.cpp
@@ -516,28 +516,6 @@ static int loadImagesOntoDevice(DeviceTy &Device) {
CurrDeviceEntryAddr = DevPtr;
}
- if (PM->getRequirements() & OMP_REQ_UNIFIED_SHARED_MEMORY ||
- PM->getRequirements() & OMPX_REQ_AUTO_ZERO_COPY) {
- AsyncInfoTy AsyncInfo(Device);
- void *DevPtr;
- Device.retrieveData(&DevPtr, CurrDeviceEntryAddr, sizeof(void *),
- AsyncInfo, /*Entry=*/nullptr, &HDTTMap);
- if (AsyncInfo.synchronize() != OFFLOAD_SUCCESS)
- return OFFLOAD_FAIL;
-
- ODBG(ODT_Mapping)
- << "Add mapping from host " << DevPtr << " to device " << DevPtr
- << " with size " << CurrDeviceEntry->Size;
-
- HDTTMap->emplace(new HostDataToTargetTy(
- (uintptr_t)DevPtr /*HstPtrBase*/,
- (uintptr_t)DevPtr /*HstPtrBegin*/,
- (uintptr_t)DevPtr + CurrHostEntry->Size /*HstPtrEnd*/,
- (uintptr_t)DevPtr /*TgtAllocBegin*/,
- (uintptr_t)DevPtr /*TgtPtrBegin*/, false /*UseHoldRefCount*/,
- nullptr, true /*IsRefCountINF*/));
- }
-
ODBG(ODT_Mapping) << "Add mapping from host " << CurrHostEntry->Address
<< " to device " << CurrDeviceEntry->Address
<< " with size " << CurrDeviceEntry->Size
diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp
index 27ee173be06f9..7c7ec49603400 100644
--- a/offload/libomptarget/omptarget.cpp
+++ b/offload/libomptarget/omptarget.cpp
@@ -988,7 +988,15 @@ postProcessingTargetDataEnd(DeviceTy *Device,
int Ret = OFFLOAD_SUCCESS;
for (auto &[HstPtrBegin, DataSize, ArgType, TPR] : EntriesInfo) {
- bool DelEntry = !TPR.isHostPointer();
+ // The reuse entry recorded on the unified-shared-memory host path has no
+ // device allocation, but it does occupy a slot in the mapping table and has
+ // to be reclaimed with the region that created it. Otherwise it lingers with
+ // a zero reference count and a later map(close, ...) of the same storage
+ // finds it and concludes the data is already on the device.
+ const bool IsHostBackedEntry =
+ TPR.getEntry() != nullptr &&
+ TPR.getEntry()->TgtAllocBegin == TPR.getEntry()->HstPtrBegin;
+ bool DelEntry = !TPR.isHostPointer() || IsHostBackedEntry;
// 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
diff --git a/offload/test/unified_shared_memory/check_tracking.c b/offload/test/unified_shared_memory/check_tracking.c
index 808c6befac52b..7607e9df90f87 100644
--- a/offload/test/unified_shared_memory/check_tracking.c
+++ b/offload/test/unified_shared_memory/check_tracking.c
@@ -14,40 +14,17 @@
#include <omp.h>
#include <stdio.h>
-#pragma omp begin declare target
-int x = 111;
-#pragma omp end declare target
-int y = 111;
-
-int present(void *p) {
- return omp_target_is_present(p, omp_get_default_device());
-}
-
int main() {
- int xl = 111;
+ int x = 111;
// CHECK: present when unmapped: 0
- printf("present when unmapped: %d\n", present(&xl));
+ printf("present when unmapped: %d\n",
+ omp_target_is_present(&x, omp_get_default_device()));
-#pragma omp target_enter_data map(alloc : xl)
+#pragma omp target_enter_data map(alloc : x)
// CHECK: present after mapping: 1
- printf("present after mapping: %d\n", present(&xl));
-#pragma omp target_exit_data map(from : xl)
- // CHECK: present after mapping: 0
- printf("present after mapping: %d\n", present(&xl));
-
- // CHECK: present when unmapped: 1 0
- printf("present when unmapped: %d %d\n", present(&x), present(&y));
-
-#pragma omp target_enter_data map(to : x, y)
-
- // CHECK: present after mapping: 1 1
- printf("present after mapping: %d %d\n", present(&x), present(&y));
-
-#pragma omp target_exit_data map(from : x, y)
-
- // CHECK: present after mapping: 1 0
- printf("present after mapping: %d %d\n", present(&x), present(&y));
+ printf("present after mapping: %d\n",
+ omp_target_is_present(&x, omp_get_default_device()));
return 0;
}
\ No newline at end of file
diff --git a/offload/test/unified_shared_memory/close.c b/offload/test/unified_shared_memory/close.c
index c813663bfbcec..83374cf4c5e67 100644
--- a/offload/test/unified_shared_memory/close.c
+++ b/offload/test/unified_shared_memory/close.c
@@ -16,6 +16,7 @@
// matching entry doesn't already exist on the device.
#include <assert.h>
+#include <stdio.h>
int x, y;
#pragma omp requires unified_shared_memory
@@ -32,4 +33,8 @@ int main() {
assert(xaddr_device == &x && "Mapped variable should not allocate on close.");
assert(yaddr_device != &y && "Unmapped variable should allocate on close.");
+
+ // CHECK: Done!
+ printf("Done!\n");
+ return 0;
}
diff --git a/offload/test/unified_shared_memory/close_data_motion.c b/offload/test/unified_shared_memory/close_data_motion.c
new file mode 100644
index 0000000000000..d5a639d444092
--- /dev/null
+++ b/offload/test/unified_shared_memory/close_data_motion.c
@@ -0,0 +1,79 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+// REQUIRES: unified_shared_memory
+// UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
+
+// amdgpu runtime crash
+// Fails on nvptx with error: an illegal memory access was encountered
+// UNSUPPORTED: amdgcn-amd-amdhsa
+// UNSUPPORTED: nvptx64-nvidia-cuda
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: intelgpu
+
+// A close map under unified shared memory gets its own device allocation, so it
+// must retain normal device data-motion semantics: values written on the device
+// have to be copied back for a `from` map, and `target update` on such a
+// mapping must actually transfer data rather than being treated as a no-op.
+//
+// This is easy to break by treating every mapping under USM as a host pointer,
+// because the copy-back and the update path are both skipped for host pointers.
+// When that happens the device writes are silently lost.
+
+#include <stdio.h>
+
+#pragma omp requires unified_shared_memory
+
+#define N 64
+
+int main() {
+ int a[N], b[N];
+
+ for (int i = 0; i < N; ++i) {
+ a[i] = 1;
+ b[i] = 1;
+ }
+
+ // The device gets its own copy of "a" because of close. The writes below must
+ // make it back to the host at the end of the region.
+#pragma omp target map(close, tofrom : a[ : N])
+ {
+ for (int i = 0; i < N; ++i)
+ a[i] += 10;
+ }
+
+ int fails = 0;
+ for (int i = 0; i < N; ++i)
+ if (a[i] != 11)
+ fails++;
+ // CHECK: close tofrom copied back: Succeeded
+ printf("close tofrom copied back: %s\n",
+ (fails == 0) ? "Succeeded" : "Failed");
+
+ // Same, but the data motion is requested explicitly with target update.
+#pragma omp target data map(close, alloc : b[ : N])
+ {
+ // Push the current host values into the device copy.
+#pragma omp target update to(b[ : N])
+
+#pragma omp target map(present, alloc : b[ : N])
+ {
+ for (int i = 0; i < N; ++i)
+ b[i] += 20;
+ }
+
+ // Pull the device values back out. If update is a no-op the host keeps 1.
+#pragma omp target update from(b[ : N])
+ }
+
+ fails = 0;
+ for (int i = 0; i < N; ++i)
+ if (b[i] != 21)
+ fails++;
+ // CHECK: close target update from: Succeeded
+ printf("close target update from: %s\n",
+ (fails == 0) ? "Succeeded" : "Failed");
+
+ // CHECK: Done!
+ printf("Done!\n");
+ return 0;
+}
diff --git a/offload/test/unified_shared_memory/close_sub_section_global.c b/offload/test/unified_shared_memory/close_sub_section_global.c
deleted file mode 100644
index 7bf128d636fa8..0000000000000
--- a/offload/test/unified_shared_memory/close_sub_section_global.c
+++ /dev/null
@@ -1,33 +0,0 @@
-// RUN: %libomptarget-compile-run-and-check-generic
-
-// REQUIRES: unified_shared_memory
-// UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
-
-// amdgpu runtime crash
-// Fails on nvptx with error: an illegal memory access was encountered
-// UNSUPPORTED: amdgcn-amd-amdhsa
-// UNSUPPORTED: nvptx64-nvidia-cuda
-// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
-// UNSUPPORTED: intelgpu
-
-#pragma omp requires unified_shared_memory
-
-#include <stdio.h>
-
-#pragma omp begin declare target
-double base[10] = {0};
-#pragma omp end declare target
-
-int main() {
-
-// close range covers base[1] and should properly
-// update the budder from the outer mapping.
-#pragma omp target map(close, alloc : base[1 : 9])
- {
- base[1] = 99.0;
- }
-#pragma omp target update from(base)
- // CHECK: base[1] = 99.000000 (expected 99.0)
- printf("base[1] = %f (expected 99.0)\n", base[1]);
- return 0;
-}
\ No newline at end of file
diff --git a/offload/test/unified_shared_memory/close_sub_section_global_oob.c b/offload/test/unified_shared_memory/close_sub_section_global_oob.c
deleted file mode 100644
index eaf5ff8c74490..0000000000000
--- a/offload/test/unified_shared_memory/close_sub_section_global_oob.c
+++ /dev/null
@@ -1,33 +0,0 @@
-// RUN: %libomptarget-compile-run-and-check-generic
-
-// REQUIRES: unified_shared_memory
-// UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
-
-// amdgpu runtime crash
-// Fails on nvptx with error: an illegal memory access was encountered
-// UNSUPPORTED: amdgcn-amd-amdhsa
-// UNSUPPORTED: nvptx64-nvidia-cuda
-// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
-// UNSUPPORTED: intelgpu
-
-#pragma omp requires unified_shared_memory
-
-#include <stdio.h>
-
-#pragma omp begin declare target
-double base[10] = {0};
-#pragma omp end declare target
-
-int main() {
-
-// close on sub-section base[2:8]; access to base[1] is outside the close
-// buffer
-#pragma omp target map(close, alloc : base[2 : 8])
- {
- base[1] = 99.0;
- }
-#pragma omp target update from(base)
- // CHECK: base[1] = 99.000000 (expected 99.0)
- printf("base[1] = %f (expected 99.0)\n", base[1]);
- return 0;
-}
\ No newline at end of file
>From 0966951002d67cd28bce66c16db50c3765bee040 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Thu, 30 Jul 2026 23:53:06 -0700
Subject: [PATCH 7/9] [OFFLOAD] Attach pointers that are newly mapped on the
USM host path
Pointer attachment is governed by a list item's mapping being new for the
construct (reference count 0 -> 1), not by device memory having been allocated for
it. Under unified shared memory a pointer mapped for the first time stays on the
host path, so it was neither recorded as newly mapped nor considered attachable,
and its deferred ATTACH was skipped.
Record such a mapping in StateInfo::NewMappings, kept apart from NewAllocations so
that the 'present' validation, which is about device allocation, is unaffected, and
let the attach lookup accept a pointer whose storage is shared with the host: that
is the storage the device dereferences, so attaching means writing the device
pointee address into it and restoring the original value at the end of the region
through the existing shadow-pointer mechanism. A host address is still rejected
for the pointee, where it would mean there is nothing to attach to.
This matters when a close mapping gives a pointee its own device allocation while
its pointer is mapped without close: the device copy of the pointer kept the host
pointee address, the kernel wrote host storage, and the close buffer was copied
back over it.
Co-Authored-By: Claude Opus 5 (1M context) <noreply at anthropic.com>
---
offload/include/OpenMP/Mapping.h | 18 ++-
offload/libomptarget/OpenMP/Mapping.cpp | 17 +--
offload/libomptarget/omptarget.cpp | 20 ++-
.../close_ptr_ptee_nested.c | 117 ++++++++++++++++++
.../close_ptr_ptee_samedir.c | 98 +++++++++++++++
5 files changed, 256 insertions(+), 14 deletions(-)
create mode 100644 offload/test/unified_shared_memory/close_ptr_ptee_nested.c
create mode 100644 offload/test/unified_shared_memory/close_ptr_ptee_samedir.c
diff --git a/offload/include/OpenMP/Mapping.h b/offload/include/OpenMP/Mapping.h
index e4024abf26690..aeed707152f00 100644
--- a/offload/include/OpenMP/Mapping.h
+++ b/offload/include/OpenMP/Mapping.h
@@ -502,10 +502,18 @@ struct StateInfoTy {
/// ATTACH map entries for deferred processing until all other maps are done.
llvm::SmallVector<AttachMapInfo> AttachEntries;
- /// Host pointers for which new memory was allocated.
+ /// Host pointers for which new device memory was allocated.
/// Key: host pointer, Value: allocation size.
+ /// Consulted by the 'present' map-type validation.
llvm::DenseMap<void *, int64_t> NewAllocations;
+ /// Host pointers whose mapping was newly created in this construct
+ /// (reference count 0 -> 1) but for which no device memory was allocated,
+ /// i.e. the unified-shared-memory host path. Kept apart from NewAllocations,
+ /// which is about device allocation.
+ /// Key: host pointer, Value: mapped size.
+ llvm::DenseMap<void *, int64_t> NewMappings;
+
/// Host pointers that had a FROM entry, but for which a data transfer was
/// skipped due to the ref-count not being zero.
/// Key: host pointer, Value: data size.
@@ -554,6 +562,14 @@ struct StateInfoTy {
return findEntryForPtr(Ptr, NewAllocations);
}
+ /// Check if a pointer's mapping was newly created in this construct, whether or
+ /// not device memory was allocated for it. Used by pointer attachment.
+ std::optional<std::pair<void *, int64_t>> wasNewlyMapped(void *Ptr) const {
+ if (auto Alloc = findEntryForPtr(Ptr, NewAllocations))
+ return Alloc;
+ return findEntryForPtr(Ptr, NewMappings);
+ }
+
/// Check if a pointer range [Ptr, Ptr+Size) is fully contained within any
/// previously completed FROM transfer.
/// Returns the matching entry if found, otherwise std::nullopt.
diff --git a/offload/libomptarget/OpenMP/Mapping.cpp b/offload/libomptarget/OpenMP/Mapping.cpp
index 295447387c607..752d5510493f0 100644
--- a/offload/libomptarget/OpenMP/Mapping.cpp
+++ b/offload/libomptarget/OpenMP/Mapping.cpp
@@ -301,13 +301,16 @@ TargetPointerResultTy MappingInfoTy::getTargetPointer(
// subsection of previously mapped allocation. The mapping would prevent
// map(close, alloc:...) from creating a new allocation as it would reuse
// the mapped allocation instead.
- LR.TPR.setEntry(
- HDTTMap
- ->emplace(new HostDataToTargetTy(
- (uintptr_t)HstPtrBase, (uintptr_t)HstPtrBegin,
- (uintptr_t)HstPtrBegin + Size, (uintptr_t)HstPtrBegin,
- (uintptr_t)HstPtrBegin, HasHoldModifier, HstPtrName))
- .first->HDTT);
+ auto Emplaced = HDTTMap->emplace(new HostDataToTargetTy(
+ (uintptr_t)HstPtrBase, (uintptr_t)HstPtrBegin,
+ (uintptr_t)HstPtrBegin + Size, (uintptr_t)HstPtrBegin,
+ (uintptr_t)HstPtrBegin, HasHoldModifier, HstPtrName));
+ LR.TPR.setEntry(Emplaced.first->HDTT);
+
+ // The mapping is new for this construct, which is what pointer attachment
+ // is governed by, so record it even though no device memory was allocated.
+ if (Emplaced.second && StateInfo)
+ StateInfo->NewMappings[HstPtrBegin] = Size;
if (Device.notifyDataMapped(HstPtrBegin, Size))
return TargetPointerResultTy{};
}
diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp
index 7c7ec49603400..ebd33384a8b0b 100644
--- a/offload/libomptarget/omptarget.cpp
+++ b/offload/libomptarget/omptarget.cpp
@@ -860,7 +860,7 @@ int processAttachEntries(DeviceTy &Device, StateInfoTy &StateInfo,
// Lambda to check if a pointer was newly allocated
auto WasNewlyAllocated = [&](void *Ptr, const char *PtrName) {
- bool WasNewlyAllocated = StateInfo.wasNewlyAllocated(Ptr).has_value();
+ bool WasNewlyAllocated = StateInfo.wasNewlyMapped(Ptr).has_value();
ODBG(ODT_Mapping) << "Attach " << PtrName << " " << Ptr
<< " was newly allocated: "
<< (WasNewlyAllocated ? "yes" : "no");
@@ -878,9 +878,15 @@ int processAttachEntries(DeviceTy &Device, StateInfoTy &StateInfo,
}
// Lambda to perform target pointer lookup and validation
+ // \p AllowHostPointer permits the lookup to succeed for storage shared with
+ // the host. That is correct for the pointer being attached: the device
+ // dereferences the same storage, so attaching writes the device pointee
+ // address into it and the original value is restored at the end of the region
+ // through the shadow-pointer mechanism. It is not correct for the pointee,
+ // where a host address would mean there is nothing to attach to.
auto LookupTargetPointer =
- [&](void *Ptr, int64_t Size,
- const char *PtrType) -> std::optional<TargetPointerResultTy> {
+ [&](void *Ptr, int64_t Size, const char *PtrType,
+ bool AllowHostPointer) -> std::optional<TargetPointerResultTy> {
// ATTACH map-type does not change ref-count, or do any allocation
// We just need to do a lookup for the pointer/pointee.
TargetPointerResultTy TPR = Device.getMappingInfo().getTgtPtrBegin(
@@ -896,7 +902,7 @@ int processAttachEntries(DeviceTy &Device, StateInfoTy &StateInfo,
<< PtrType << " not present on device";
return std::nullopt;
}
- if (TPR.Flags.IsHostPointer) {
+ if (TPR.Flags.IsHostPointer && !AllowHostPointer) {
ODBG(ODT_Mapping) << "Skipping ATTACH entry " << EntryIdx
<< ": device version of the " << PtrType
<< " is a host pointer.";
@@ -909,7 +915,8 @@ int processAttachEntries(DeviceTy &Device, StateInfoTy &StateInfo,
// Get device version of the pointee (e.g., &p[10]) first, as we can
// release its TPR after extracting the pointer value.
void *TgtPteeBegin = [&]() -> void * {
- if (auto PteeTPROpt = LookupTargetPointer(HstPteeBegin, 0, "pointee"))
+ if (auto PteeTPROpt = LookupTargetPointer(HstPteeBegin, 0, "pointee",
+ /*AllowHostPointer=*/false))
return PteeTPROpt->TargetPointer;
return nullptr;
}();
@@ -919,7 +926,8 @@ int processAttachEntries(DeviceTy &Device, StateInfoTy &StateInfo,
// Get device version of the pointer (e.g., &p) next. We need to keep its
// TPR for use in shadow-pointer handling during pointer-attachment.
- auto PtrTPROpt = LookupTargetPointer(HstPtr, PtrSize, "pointer");
+ auto PtrTPROpt = LookupTargetPointer(HstPtr, PtrSize, "pointer",
+ /*AllowHostPointer=*/true);
if (!PtrTPROpt)
continue;
TargetPointerResultTy &PtrTPR = *PtrTPROpt;
diff --git a/offload/test/unified_shared_memory/close_ptr_ptee_nested.c b/offload/test/unified_shared_memory/close_ptr_ptee_nested.c
new file mode 100644
index 0000000000000..c135ef3a7033b
--- /dev/null
+++ b/offload/test/unified_shared_memory/close_ptr_ptee_nested.c
@@ -0,0 +1,117 @@
+// Pointer + pointee mapping with the `close` map-type modifier under unified
+// shared memory, where the pointer `p` and its pointee `p[0:10]` are mapped on
+// SEPARATE, nested `target data` directives.
+//
+// Under USM the `close` modifier forces a real device allocation for the list
+// item even though the memory is shared. Whenever the pointer is (newly) mapped
+// alongside a device-resident pointee, it must be pointer-attached to that
+// device pointee, and its original host value must be restored at the end of
+// the region via shadow-pointer tracking -- otherwise the host pointer would be
+// left pointing at the device allocation, and (for a close pointee) the device
+// buffer would diverge from the host storage the kernel writes through the
+// unattached host pointer.
+//
+// Each variation is run twice: once under LIBOMPTARGET_DEBUG=1 to check the
+// pointer-attachment bookkeeping, and once to check the program results.
+//
+// RUN: %libomptarget-compile-generic -DV1=1
+// RUN: env LIBOMPTARGET_DEBUG=1 %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefixes=ATTACHED
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefixes=V1,ALL
+//
+// RUN: %libomptarget-compile-generic -DV2=1
+// RUN: env LIBOMPTARGET_DEBUG=1 %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefixes=ATTACHED
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefixes=V2,ALL
+//
+// RUN: %libomptarget-compile-generic
+// RUN: env LIBOMPTARGET_DEBUG=1 %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefixes=ATTACHED
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefixes=V3,ALL
+//
+// REQUIRES: unified_shared_memory
+// REQUIRES: libomptarget-debug
+//
+// UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
+// UNSUPPORTED: amdgcn-amd-amdhsa
+// UNSUPPORTED: nvptx64-nvidia-cuda
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: intelgpu
+
+#include <omp.h>
+#include <stdio.h>
+
+#pragma omp requires unified_shared_memory
+
+int x[10] = {0};
+int *p = &x[0];
+int *p_device = NULL;
+int **paddr_device = NULL;
+
+int main() {
+ printf("Before tgt: p %s p_host\n", p == &x[0] ? "==" : "!=");
+
+#if V1
+#pragma omp target data map(tofrom : p)
+#elif V2
+#pragma omp target data map(close, tofrom : p)
+#else // V3
+#pragma omp target data map(close, tofrom : p[0 : 10])
+#endif
+ {
+#if V1
+#pragma omp target data map(close, tofrom : p[0 : 10])
+#elif V2
+#pragma omp target data map(tofrom : p[0 : 10])
+#else // V3
+#pragma omp target data map(tofrom : p) map(tofrom : p[0 : 10])
+#endif
+ {
+#pragma omp target map(present, alloc : p) map(from : p_device, paddr_device)
+ {
+ p_device = p;
+ paddr_device = &p;
+ p[0] = 111;
+ }
+ }
+ }
+
+ printf("In tgt: p_device %s p_host\n", p_device == &x[0] ? "==" : "!=");
+ printf("In tgt: paddr_device %s &p_host\n", paddr_device == &p ? "==" : "!=");
+ printf("After tgt: p %s p_host\n", p == &x[0] ? "==" : "!=");
+ printf("p[0] = %d\n", x[0]);
+
+ // When the pointee has a device allocation of its own, the pointer is attached
+ // to it and its original host value is restored at the end of the region.
+ // ATTACHED: ATTACH entry {{.*}} processed successfully
+ // ATTACHED: Restoring host pointer
+
+
+ // ALL: Before tgt: p == p_host
+
+ // The pointer is attached to the close-allocated device pointee, so its
+ // device value differs from the host address.
+ // V1: In tgt: p_device != p_host
+ // V1: In tgt: paddr_device == &p_host
+
+ // p itself is close-allocated, giving the pointer variable its own device
+ // storage (so &p differs on the device); the pointee stays on the USM host
+ // path, so the device pointer value equals the host address.
+ // V2: In tgt: p_device == p_host
+ // V2: In tgt: paddr_device != &p_host
+
+ // The close pointee (mapped by the enclosing region) has a device
+ // allocation, and p -- newly mapped on the USM host path by the inner region
+ // -- is attached to it, so its device value differs from the host address
+ // while the pointer variable itself stays on the host path (&p matches).
+ // V3: In tgt: p_device != p_host
+ // V3: In tgt: paddr_device == &p_host
+
+ // The host pointer must be intact afterwards and the kernel's write must be
+ // visible on the host.
+ // ALL: After tgt: p == p_host
+ // ALL: p[0] = 111
+}
diff --git a/offload/test/unified_shared_memory/close_ptr_ptee_samedir.c b/offload/test/unified_shared_memory/close_ptr_ptee_samedir.c
new file mode 100644
index 0000000000000..ae8cb7985881a
--- /dev/null
+++ b/offload/test/unified_shared_memory/close_ptr_ptee_samedir.c
@@ -0,0 +1,98 @@
+// Pointer + pointee mapping with the `close` map-type modifier under unified
+// shared memory, where the pointer `p` and its pointee `p[0:10]` are mapped on
+// the SAME `target data` directive.
+//
+// Companion to close_ptr_ptee_nested.c; see that file for the rationale. Here
+// the pointer and pointee are processed together in a single targetDataBegin,
+// so pointer-attachment sees both the pointer and the (close-allocated) pointee
+// in the same construct.
+//
+// Each variation is run twice: once under LIBOMPTARGET_DEBUG=1 to check the
+// pointer-attachment bookkeeping, and once to check the program results.
+//
+// RUN: %libomptarget-compile-generic -DV1=1
+// RUN: env LIBOMPTARGET_DEBUG=1 %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefixes=ATTACHED
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefixes=V1,ALL
+//
+// RUN: %libomptarget-compile-generic -DV2=1
+// RUN: env LIBOMPTARGET_DEBUG=1 %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefixes=ATTACHED
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefixes=V2,ALL
+//
+// RUN: %libomptarget-compile-generic
+// RUN: env LIBOMPTARGET_DEBUG=1 %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefixes=ATTACHED
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefixes=V3,ALL
+//
+// REQUIRES: unified_shared_memory
+// REQUIRES: libomptarget-debug
+//
+// UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
+// UNSUPPORTED: amdgcn-amd-amdhsa
+// UNSUPPORTED: nvptx64-nvidia-cuda
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: intelgpu
+
+#include <omp.h>
+#include <stdio.h>
+
+#pragma omp requires unified_shared_memory
+
+int x[10] = {0};
+int *p = &x[0];
+int *p_device = NULL;
+int **paddr_device = NULL;
+
+int main() {
+ printf("Before tgt: p %s p_host\n", p == &x[0] ? "==" : "!=");
+
+#if V1
+#pragma omp target data map(tofrom : p) map(close, tofrom : p[0 : 10])
+#elif V2
+#pragma omp target data map(close, tofrom : p) map(tofrom : p[0 : 10])
+#else // V3
+#pragma omp target data map(close, tofrom : p) map(close, tofrom : p[0 : 10])
+#endif
+ {
+#pragma omp target map(present, alloc : p) map(from : p_device, paddr_device)
+ {
+ p_device = p;
+ paddr_device = &p;
+ p[0] = 111;
+ }
+ }
+
+ printf("In tgt: p_device %s p_host\n", p_device == &x[0] ? "==" : "!=");
+ printf("In tgt: paddr_device %s &p_host\n", paddr_device == &p ? "==" : "!=");
+ printf("After tgt: p %s p_host\n", p == &x[0] ? "==" : "!=");
+ printf("p[0] = %d\n", x[0]);
+
+ // When the pointee has a device allocation of its own, the pointer is attached
+ // to it and its original host value is restored at the end of the region.
+ // ATTACHED: ATTACH entry {{.*}} processed successfully
+ // ATTACHED: Restoring host pointer
+
+
+ // ALL: Before tgt: p == p_host
+
+ // The close pointee is device-allocated and p (host path) is attached to it.
+ // V1: In tgt: p_device != p_host
+ // V1: In tgt: paddr_device == &p_host
+
+ // p is close-allocated (own device storage, so &p differs); the pointee is a
+ // plain USM mapping, so the device pointer value equals the host address.
+ // V2: In tgt: p_device == p_host
+ // V2: In tgt: paddr_device != &p_host
+
+ // Both p and the pointee are close-allocated, so both the device address of p
+ // and its attached value differ from the host.
+ // V3: In tgt: p_device != p_host
+ // V3: In tgt: paddr_device != &p_host
+
+ // ALL: After tgt: p == p_host
+ // ALL: p[0] = 111
+}
>From a2adef1ecfac71c811e965f36f74113456bce4dc Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Thu, 30 Jul 2026 23:59:52 -0700
Subject: [PATCH 8/9] [OpenMP] Communicate declare-target variable size under
unified shared memory
Under unified shared memory a declare-target variable is represented on the device
by a reference pointer to the host storage, and the offload entry emitted for it
describes that pointer:
no USM: { Address = &arr, Size = sizeof(arr) }
with USM: { Address = &arr_decl_tgt_ref_ptr, Size = sizeof(void *) }
Address and Size have to keep describing the pointer, since that is what
get_global() looks up and data_submit() writes, so the size of the variable itself
is communicated in addition, in the entry's spare Data field. The variable's
address does not need to be sent: the host copy of the reference pointer holds it,
so the runtime obtains it with a plain dereference.
With that, the runtime registers the variable's storage with its real extent, and:
- omp_target_is_present() reports declare-target storage as present, for its
whole extent rather than its first sizeof(void *) bytes.
- A close mapping of such storage is recognized as overlapping it, so it reuses
the existing device buffer instead of getting its own. That fixes the lost
write when the mapping covers part of the object or is reached through a
pointer into it, and lets close_sub_section_global{,_oob}.c pass, the latter of
which used to abort.
Old and new components interoperate: the size only ever occupies a field that was
previously zero, and the runtime acts on it only when it is set. An old binary on a
new runtime behaves as it did before, and a new binary on an old runtime has the
size ignored.
Flang needs no separate change, reaching the same
OpenMPIRBuilder::registerTargetGlobalVariable() through OpenMPToLLVMIRTranslation;
the size is derived from the LLVM module rather than from anything
language-specific.
Co-Authored-By: Claude Opus 5 (1M context) <noreply at anthropic.com>
---
.../declare_target_local_usm_codegen.cpp | 2 +-
..._target_requires_unified_shared_memory.cpp | 4 +-
.../llvm/Frontend/OpenMP/OMPIRBuilder.h | 17 ++++-
llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 33 +++++++--
offload/include/OpenMP/Mapping.h | 4 +-
offload/libomptarget/OpenMP/Mapping.cpp | 3 +-
offload/libomptarget/PluginManager.cpp | 30 ++++++++
offload/libomptarget/omptarget.cpp | 12 +--
.../unified_shared_memory/check_tracking.c | 38 ++++++++--
.../close_alloc_declare_target_lifetime.c | 74 +++++++++++++++++++
.../close_ptr_ptee_nested.c | 9 +--
.../close_ptr_ptee_samedir.c | 9 +--
.../close_sub_section_global.c | 33 +++++++++
.../close_sub_section_global_oob.c | 33 +++++++++
.../declare_target_map_extent.c | 67 +++++++++++++++++
.../declare_target_tracking.c | 65 ++++++++++++++++
16 files changed, 397 insertions(+), 36 deletions(-)
create mode 100644 offload/test/unified_shared_memory/close_alloc_declare_target_lifetime.c
create mode 100644 offload/test/unified_shared_memory/close_sub_section_global.c
create mode 100644 offload/test/unified_shared_memory/close_sub_section_global_oob.c
create mode 100644 offload/test/unified_shared_memory/declare_target_map_extent.c
create mode 100644 offload/test/unified_shared_memory/declare_target_tracking.c
diff --git a/clang/test/OpenMP/declare_target_local_usm_codegen.cpp b/clang/test/OpenMP/declare_target_local_usm_codegen.cpp
index d97d6f409d265..2b199f1d4aaa3 100644
--- a/clang/test/OpenMP/declare_target_local_usm_codegen.cpp
+++ b/clang/test/OpenMP/declare_target_local_usm_codegen.cpp
@@ -31,7 +31,7 @@ int enter_var;
// enter_var with USM: pointer-reference indirection
// HOST-DAG: @enter_var_decl_tgt_ref_ptr = weak global ptr @enter_var
-// HOST-DAG: @.offloading.entry.enter_var_decl_tgt_ref_ptr = weak constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 0, ptr @enter_var_decl_tgt_ref_ptr, ptr @.offloading.entry_name{{.*}}, i64 8, i64 0, ptr null }, section "llvm_offload_entries"
+// HOST-DAG: @.offloading.entry.enter_var_decl_tgt_ref_ptr = weak constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 0, ptr @enter_var_decl_tgt_ref_ptr, ptr @.offloading.entry_name{{.*}}, i64 8, i64 4, ptr null }, section "llvm_offload_entries"
// Device: local_var is a direct global, enter_var uses ref ptr
// DEVICE-DAG: @local_var = protected addrspace(1) global i32 0
diff --git a/clang/test/OpenMP/nvptx_target_requires_unified_shared_memory.cpp b/clang/test/OpenMP/nvptx_target_requires_unified_shared_memory.cpp
index b39007a712994..e0e4c1f1ea66b 100644
--- a/clang/test/OpenMP/nvptx_target_requires_unified_shared_memory.cpp
+++ b/clang/test/OpenMP/nvptx_target_requires_unified_shared_memory.cpp
@@ -47,10 +47,10 @@ int bar(int n){
// CHECK-HOST: [[OFFLOAD_MAPTYPES:@.+]] = private unnamed_addr constant [3 x i64] [i64 800, i64 800, i64 288]
// CHECK-HOST: [[OMP_OFFLOAD_ENTRY_LINK_VAR_PTR_NAME:@.+]] = internal unnamed_addr constant [21 x i8]
-// CHECK-HOST: [[OMP_OFFLOAD_ENTRY_LINK_VAR_PTR:@.+]] = weak constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 1, ptr [[VAR_DECL_TGT_LINK_PTR]], ptr [[OMP_OFFLOAD_ENTRY_LINK_VAR_PTR_NAME]], i64 8, i64 0, ptr null }, section "llvm_offload_entries"
+// CHECK-HOST: [[OMP_OFFLOAD_ENTRY_LINK_VAR_PTR:@.+]] = weak constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 1, ptr [[VAR_DECL_TGT_LINK_PTR]], ptr [[OMP_OFFLOAD_ENTRY_LINK_VAR_PTR_NAME]], i64 8, i64 8, ptr null }, section "llvm_offload_entries"
// CHECK-HOST: [[OMP_OFFLOAD_ENTRY_TO_VAR_PTR_NAME:@.+]] = internal unnamed_addr constant [24 x i8]
-// CHECK-HOST: [[OMP_OFFLOAD_ENTRY_TO_VAR_PTR:@.+]] = weak constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 0, ptr [[VAR_DECL_TGT_TO_PTR]], ptr @.offloading.entry_name.1, i64 8, i64 0, ptr null }, section "llvm_offload_entries"
+// CHECK-HOST: [[OMP_OFFLOAD_ENTRY_TO_VAR_PTR:@.+]] = weak constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 0, ptr [[VAR_DECL_TGT_TO_PTR]], ptr @.offloading.entry_name.1, i64 8, i64 8, ptr null }, section "llvm_offload_entries"
// CHECK-HOST: [[N_CASTED:%.+]] = alloca i64
// CHECK-HOST: [[SUM_CASTED:%.+]] = alloca i64
diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
index 965ac358c259e..2277e28e5cd1f 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
+++ b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
@@ -423,6 +423,9 @@ class OffloadEntriesInfoManager {
class OffloadEntryInfoDeviceGlobalVar final : public OffloadEntryInfo {
/// Type of the global variable.
int64_t VarSize;
+ /// Size of the variable a device reference pointer refers to, or 0 when the
+ /// entry describes the variable itself. See createOffloadEntry().
+ int64_t PointeeSize = 0;
GlobalValue::LinkageTypes Linkage;
const std::string VarName;
@@ -445,6 +448,8 @@ class OffloadEntriesInfoManager {
int64_t getVarSize() const { return VarSize; }
StringRef getVarName() const { return VarName; }
void setVarSize(int64_t Size) { VarSize = Size; }
+ int64_t getPointeeSize() const { return PointeeSize; }
+ void setPointeeSize(int64_t Size) { PointeeSize = Size; }
GlobalValue::LinkageTypes getLinkage() const { return Linkage; }
void setLinkage(GlobalValue::LinkageTypes LT) { Linkage = LT; }
static bool classof(const OffloadEntryInfo *Info) {
@@ -460,7 +465,8 @@ class OffloadEntriesInfoManager {
/// Register device global variable entry.
LLVM_ABI void registerDeviceGlobalVarEntryInfo(
StringRef VarName, Constant *Addr, int64_t VarSize,
- OMPTargetGlobalVarEntryKind Flags, GlobalValue::LinkageTypes Linkage);
+ OMPTargetGlobalVarEntryKind Flags, GlobalValue::LinkageTypes Linkage,
+ int64_t PointeeSize = 0);
/// Checks if the variable with the given name has been registered already.
bool hasDeviceGlobalVarEntryInfo(StringRef VarName) const {
return OffloadEntriesDeviceGlobalVar.count(VarName) > 0;
@@ -3014,9 +3020,16 @@ class OpenMPIRBuilder {
/// Creates offloading entry for the provided entry ID \a ID, address \a
/// Addr, size \a Size, and flags \a Flags.
+ ///
+ /// \a PointeeSize is nonzero only for a declare-target variable under unified
+ /// shared memory, where \a Addr is a device reference pointer to the host
+ /// variable and \a Size is therefore the size of that pointer. It carries the
+ /// size of the variable itself, which the runtime needs in order to register
+ /// the variable's storage rather than the pointer's.
LLVM_ABI void createOffloadEntry(Constant *ID, Constant *Addr, uint64_t Size,
int32_t Flags, GlobalValue::LinkageTypes,
- StringRef Name = "");
+ StringRef Name = "",
+ uint64_t PointeeSize = 0);
/// The kind of errors that can occur when emitting the offload entries and
/// metadata.
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index f8c1999fe1b89..f58e45779a8e4 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -11874,11 +11874,12 @@ std::unique_ptr<CodeExtractor> DeviceSharedMemOutlineInfo::createCodeExtractor(
void OpenMPIRBuilder::createOffloadEntry(Constant *ID, Constant *Addr,
uint64_t Size, int32_t Flags,
GlobalValue::LinkageTypes,
- StringRef Name) {
+ StringRef Name, uint64_t PointeeSize) {
if (!Config.isGPU()) {
llvm::offloading::emitOffloadingEntry(
M, object::OffloadKind::OFK_OpenMP, ID,
- Name.empty() ? Addr->getName() : Name, Size, Flags, /*Data=*/0);
+ Name.empty() ? Addr->getName() : Name, Size, Flags,
+ /*Data=*/PointeeSize);
return;
}
// TODO: Add support for global variables on the device after declare target
@@ -12051,7 +12052,8 @@ void OpenMPIRBuilder::createOffloadEntriesAndInfoMetadata(
Flags, CE->getLinkage(), CE->getVarName());
else
createOffloadEntry(CE->getAddress(), CE->getAddress(), CE->getVarSize(),
- Flags, CE->getLinkage());
+ Flags, CE->getLinkage(), /*Name=*/"",
+ CE->getPointeeSize());
} else {
llvm_unreachable("Unsupported entry kind.");
@@ -12224,6 +12226,7 @@ void OpenMPIRBuilder::registerTargetGlobalVariable(
OffloadEntriesInfoManager::OMPTargetGlobalVarEntryKind Flags;
StringRef VarName;
int64_t VarSize;
+ int64_t PointeeSize = 0;
GlobalValue::LinkageTypes Linkage;
if ((CaptureClause == OffloadEntriesInfoManager::OMPTargetGlobalVarEntryTo ||
@@ -12280,10 +12283,25 @@ void OpenMPIRBuilder::registerTargetGlobalVariable(
}
VarSize = M.getDataLayout().getPointerSize();
Linkage = GlobalValue::WeakAnyLinkage;
+
+ // Under unified shared memory Addr is a device reference pointer to the
+ // host variable, so VarSize above is the size of that pointer. Record the
+ // size of the variable itself as well: the runtime needs it to register the
+ // variable's storage, which is what a map of the variable refers to.
+ //
+ // This only applies to storage that lives on the host for the whole
+ // program. A "declare target link" variable without unified shared memory
+ // is also reached through a reference pointer, but its storage is on the
+ // device, so there is nothing for the runtime to register here.
+ if (Config.hasRequiresUnifiedSharedMemory() && !IsDeclaration &&
+ !Config.isTargetDevice())
+ if (GlobalValue *LlvmVal = M.getNamedValue(MangledName))
+ PointeeSize = divideCeil(
+ M.getDataLayout().getTypeSizeInBits(LlvmVal->getValueType()), 8);
}
- OffloadInfoManager.registerDeviceGlobalVarEntryInfo(VarName, Addr, VarSize,
- Flags, Linkage);
+ OffloadInfoManager.registerDeviceGlobalVarEntryInfo(
+ VarName, Addr, VarSize, Flags, Linkage, PointeeSize);
}
/// Loads all the offload entries information from the host IR
@@ -12728,7 +12746,8 @@ void OffloadEntriesInfoManager::initializeDeviceGlobalVarEntryInfo(
void OffloadEntriesInfoManager::registerDeviceGlobalVarEntryInfo(
StringRef VarName, Constant *Addr, int64_t VarSize,
- OMPTargetGlobalVarEntryKind Flags, GlobalValue::LinkageTypes Linkage) {
+ OMPTargetGlobalVarEntryKind Flags, GlobalValue::LinkageTypes Linkage,
+ int64_t PointeeSize) {
if (OMPBuilder->Config.isTargetDevice()) {
// This could happen if the device compilation is invoked standalone.
if (!hasDeviceGlobalVarEntryInfo(VarName))
@@ -12742,6 +12761,7 @@ void OffloadEntriesInfoManager::registerDeviceGlobalVarEntryInfo(
return;
}
Entry.setVarSize(VarSize);
+ Entry.setPointeeSize(PointeeSize);
Entry.setLinkage(Linkage);
Entry.setAddress(Addr);
} else {
@@ -12764,6 +12784,7 @@ void OffloadEntriesInfoManager::registerDeviceGlobalVarEntryInfo(
else
OffloadEntriesDeviceGlobalVar.try_emplace(
VarName, OffloadingEntriesNum, Addr, VarSize, Flags, Linkage, "");
+ OffloadEntriesDeviceGlobalVar[VarName].setPointeeSize(PointeeSize);
++OffloadingEntriesNum;
}
}
diff --git a/offload/include/OpenMP/Mapping.h b/offload/include/OpenMP/Mapping.h
index aeed707152f00..4f20cc5490143 100644
--- a/offload/include/OpenMP/Mapping.h
+++ b/offload/include/OpenMP/Mapping.h
@@ -562,8 +562,8 @@ struct StateInfoTy {
return findEntryForPtr(Ptr, NewAllocations);
}
- /// Check if a pointer's mapping was newly created in this construct, whether or
- /// not device memory was allocated for it. Used by pointer attachment.
+ /// Check if a pointer's mapping was newly created in this construct, whether
+ /// or not device memory was allocated for it. Used by pointer attachment.
std::optional<std::pair<void *, int64_t>> wasNewlyMapped(void *Ptr) const {
if (auto Alloc = findEntryForPtr(Ptr, NewAllocations))
return Alloc;
diff --git a/offload/libomptarget/OpenMP/Mapping.cpp b/offload/libomptarget/OpenMP/Mapping.cpp
index 752d5510493f0..dbf88118000fa 100644
--- a/offload/libomptarget/OpenMP/Mapping.cpp
+++ b/offload/libomptarget/OpenMP/Mapping.cpp
@@ -308,7 +308,8 @@ TargetPointerResultTy MappingInfoTy::getTargetPointer(
LR.TPR.setEntry(Emplaced.first->HDTT);
// The mapping is new for this construct, which is what pointer attachment
- // is governed by, so record it even though no device memory was allocated.
+ // is governed by, so record it even though no device memory was
+ // allocated.
if (Emplaced.second && StateInfo)
StateInfo->NewMappings[HstPtrBegin] = Size;
if (Device.notifyDataMapped(HstPtrBegin, Size))
diff --git a/offload/libomptarget/PluginManager.cpp b/offload/libomptarget/PluginManager.cpp
index 41b653a60adfd..3d3befe095a3f 100644
--- a/offload/libomptarget/PluginManager.cpp
+++ b/offload/libomptarget/PluginManager.cpp
@@ -516,6 +516,36 @@ static int loadImagesOntoDevice(DeviceTy &Device) {
CurrDeviceEntryAddr = DevPtr;
}
+ // Under unified shared memory a declare-target variable is represented
+ // on the device by a reference pointer to the host storage, so the
+ // entry below describes that pointer. Code generation records the size
+ // of the variable itself in Data; register its storage as well, so that
+ // it is reported as present and a mapping of it, or of part of it, is
+ // recognized as referring to storage already on the device.
+ //
+ // The host copy of the reference pointer holds the variable's address,
+ // so no device read is needed to find it.
+ if ((PM->getRequirements() & OMP_REQ_UNIFIED_SHARED_MEMORY ||
+ PM->getRequirements() & OMPX_REQ_AUTO_ZERO_COPY) &&
+ CurrHostEntry->Data) {
+ void *HostVar = *reinterpret_cast<void **>(CurrHostEntry->Address);
+ uint64_t VarSize = CurrHostEntry->Data;
+
+ ODBG(ODT_Mapping)
+ << "Add mapping from host " << HostVar << " to device " << HostVar
+ << " with size " << VarSize << ", name \""
+ << CurrDeviceEntry->SymbolName << "\" (declare-target storage)";
+
+ HDTTMap->emplace(new HostDataToTargetTy(
+ (uintptr_t)HostVar, (uintptr_t)HostVar,
+ (uintptr_t)HostVar + VarSize, (uintptr_t)HostVar,
+ (uintptr_t)HostVar, false /*UseHoldRefCount*/,
+ CurrHostEntry->SymbolName, true /*IsRefCountINF*/));
+
+ if (Device.notifyDataMapped(HostVar, VarSize))
+ return OFFLOAD_FAIL;
+ }
+
ODBG(ODT_Mapping) << "Add mapping from host " << CurrHostEntry->Address
<< " to device " << CurrDeviceEntry->Address
<< " with size " << CurrDeviceEntry->Size
diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp
index ebd33384a8b0b..917b865da53f5 100644
--- a/offload/libomptarget/omptarget.cpp
+++ b/offload/libomptarget/omptarget.cpp
@@ -881,9 +881,9 @@ int processAttachEntries(DeviceTy &Device, StateInfoTy &StateInfo,
// \p AllowHostPointer permits the lookup to succeed for storage shared with
// the host. That is correct for the pointer being attached: the device
// dereferences the same storage, so attaching writes the device pointee
- // address into it and the original value is restored at the end of the region
- // through the shadow-pointer mechanism. It is not correct for the pointee,
- // where a host address would mean there is nothing to attach to.
+ // address into it and the original value is restored at the end of the
+ // region through the shadow-pointer mechanism. It is not correct for the
+ // pointee, where a host address would mean there is nothing to attach to.
auto LookupTargetPointer =
[&](void *Ptr, int64_t Size, const char *PtrType,
bool AllowHostPointer) -> std::optional<TargetPointerResultTy> {
@@ -998,9 +998,9 @@ postProcessingTargetDataEnd(DeviceTy *Device,
for (auto &[HstPtrBegin, DataSize, ArgType, TPR] : EntriesInfo) {
// The reuse entry recorded on the unified-shared-memory host path has no
// device allocation, but it does occupy a slot in the mapping table and has
- // to be reclaimed with the region that created it. Otherwise it lingers with
- // a zero reference count and a later map(close, ...) of the same storage
- // finds it and concludes the data is already on the device.
+ // to be reclaimed with the region that created it. Otherwise it lingers
+ // with a zero reference count and a later map(close, ...) of the same
+ // storage finds it and concludes the data is already on the device.
const bool IsHostBackedEntry =
TPR.getEntry() != nullptr &&
TPR.getEntry()->TgtAllocBegin == TPR.getEntry()->HstPtrBegin;
diff --git a/offload/test/unified_shared_memory/check_tracking.c b/offload/test/unified_shared_memory/check_tracking.c
index 7607e9df90f87..8b9c1964578cb 100644
--- a/offload/test/unified_shared_memory/check_tracking.c
+++ b/offload/test/unified_shared_memory/check_tracking.c
@@ -14,17 +14,43 @@
#include <omp.h>
#include <stdio.h>
+#pragma omp begin declare target
+int x = 111;
+#pragma omp end declare target
+int y = 111;
+
+int present(void *p) {
+ return omp_target_is_present(p, omp_get_default_device());
+}
+
int main() {
- int x = 111;
+ int xl = 111;
// CHECK: present when unmapped: 0
- printf("present when unmapped: %d\n",
- omp_target_is_present(&x, omp_get_default_device()));
+ printf("present when unmapped: %d\n", present(&xl));
-#pragma omp target_enter_data map(alloc : x)
+#pragma omp target_enter_data map(alloc : xl)
// CHECK: present after mapping: 1
- printf("present after mapping: %d\n",
- omp_target_is_present(&x, omp_get_default_device()));
+ printf("present after mapping: %d\n", present(&xl));
+#pragma omp target_exit_data map(from : xl)
+ // CHECK: present after mapping: 0
+ printf("present after mapping: %d\n", present(&xl));
+
+ // x is declare-target, so it is in the device data environment from the
+ // start, whereas y is an ordinary global and is not mapped yet.
+ // CHECK: present when unmapped: 1 0
+ printf("present when unmapped: %d %d\n", present(&x), present(&y));
+
+#pragma omp target_enter_data map(to : x, y)
+
+ // CHECK: present after mapping: 1 1
+ printf("present after mapping: %d %d\n", present(&x), present(&y));
+
+#pragma omp target_exit_data map(from : x, y)
+
+ // x stays present after the exit, being declare-target; y does not.
+ // CHECK: present after mapping: 1 0
+ printf("present after mapping: %d %d\n", present(&x), present(&y));
return 0;
}
\ No newline at end of file
diff --git a/offload/test/unified_shared_memory/close_alloc_declare_target_lifetime.c b/offload/test/unified_shared_memory/close_alloc_declare_target_lifetime.c
new file mode 100644
index 0000000000000..80bd58d8acff0
--- /dev/null
+++ b/offload/test/unified_shared_memory/close_alloc_declare_target_lifetime.c
@@ -0,0 +1,74 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+// REQUIRES: unified_shared_memory
+// UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
+
+// amdgpu runtime crash
+// Fails on nvptx with error: an illegal memory access was encountered
+// UNSUPPORTED: amdgcn-amd-amdhsa
+// UNSUPPORTED: nvptx64-nvidia-cuda
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: intelgpu
+
+// A declare-target variable is in the device data environment for the whole
+// program, so its storage outlives any individual mapping of it, and a
+// `target update` on it after some unrelated region has ended is valid user
+// code.
+//
+// This test documents that the implementation-specific optimization of giving a
+// `close` mapping its own device allocation conflicts with that: with
+// map(close, alloc : ...) there is no copy-back and the separate allocation is
+// released at the end of the region, so a value the device wrote there is lost.
+//
+// The optimization applies only to storage that is not already on the device. A
+// mapping that lies within such storage stays on the host path and so shares
+// its device buffer, which is what the containment check in getTargetPointer()
+// expresses. For a declare-target variable that requires knowing the variable's
+// extent, which code generation communicates in its offload entry: under
+// unified shared memory the entry otherwise describes only the device reference
+// pointer.
+
+#include <stdio.h>
+
+#pragma omp requires unified_shared_memory
+
+#pragma omp begin declare target
+double part[10] = {0};
+double viaptr[10] = {0};
+#pragma omp end declare target
+
+// An unrelated variable, mapped alongside the close mapping below. Its entry is
+// what makes the lookup for the close mapping find a neighbor, so that the
+// mapping is not left on the host path.
+double *probe;
+
+int main() {
+ // (A) A close mapping of part of a declare-target object.
+ //
+ // CHECK: A: 11.000000
+#pragma omp target map(close, alloc : part[3 : 4]) map(from : probe)
+ {
+ probe = &part[3];
+ part[3] = 11.0;
+ }
+ // Legal: "part" is still in the device data environment.
+#pragma omp target update from(part[3 : 1])
+ printf("A: %f\n", part[3]);
+
+ // (B) The same, reached through a pointer into the middle of the object, so
+ // the mapping carries no reference to the declare-target variable at all.
+ //
+ // CHECK: B: 22.000000
+ double *p = &viaptr[6];
+#pragma omp target map(close, alloc : p[0 : 3]) map(from : probe)
+ {
+ probe = &p[0];
+ p[0] = 22.0;
+ }
+#pragma omp target update from(viaptr[6 : 1])
+ printf("B: %f\n", viaptr[6]);
+
+ // CHECK: Done!
+ printf("Done!\n");
+ return 0;
+}
diff --git a/offload/test/unified_shared_memory/close_ptr_ptee_nested.c b/offload/test/unified_shared_memory/close_ptr_ptee_nested.c
index c135ef3a7033b..62d363efa8ed8 100644
--- a/offload/test/unified_shared_memory/close_ptr_ptee_nested.c
+++ b/offload/test/unified_shared_memory/close_ptr_ptee_nested.c
@@ -84,11 +84,10 @@ int main() {
printf("After tgt: p %s p_host\n", p == &x[0] ? "==" : "!=");
printf("p[0] = %d\n", x[0]);
- // When the pointee has a device allocation of its own, the pointer is attached
- // to it and its original host value is restored at the end of the region.
- // ATTACHED: ATTACH entry {{.*}} processed successfully
- // ATTACHED: Restoring host pointer
-
+ // When the pointee has a device allocation of its own, the pointer is
+ // attached to it and its original host value is restored at the end of the
+ // region. ATTACHED: ATTACH entry {{.*}} processed successfully ATTACHED:
+ // Restoring host pointer
// ALL: Before tgt: p == p_host
diff --git a/offload/test/unified_shared_memory/close_ptr_ptee_samedir.c b/offload/test/unified_shared_memory/close_ptr_ptee_samedir.c
index ae8cb7985881a..78bdc326e1da3 100644
--- a/offload/test/unified_shared_memory/close_ptr_ptee_samedir.c
+++ b/offload/test/unified_shared_memory/close_ptr_ptee_samedir.c
@@ -71,11 +71,10 @@ int main() {
printf("After tgt: p %s p_host\n", p == &x[0] ? "==" : "!=");
printf("p[0] = %d\n", x[0]);
- // When the pointee has a device allocation of its own, the pointer is attached
- // to it and its original host value is restored at the end of the region.
- // ATTACHED: ATTACH entry {{.*}} processed successfully
- // ATTACHED: Restoring host pointer
-
+ // When the pointee has a device allocation of its own, the pointer is
+ // attached to it and its original host value is restored at the end of the
+ // region. ATTACHED: ATTACH entry {{.*}} processed successfully ATTACHED:
+ // Restoring host pointer
// ALL: Before tgt: p == p_host
diff --git a/offload/test/unified_shared_memory/close_sub_section_global.c b/offload/test/unified_shared_memory/close_sub_section_global.c
new file mode 100644
index 0000000000000..7bf128d636fa8
--- /dev/null
+++ b/offload/test/unified_shared_memory/close_sub_section_global.c
@@ -0,0 +1,33 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+// REQUIRES: unified_shared_memory
+// UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
+
+// amdgpu runtime crash
+// Fails on nvptx with error: an illegal memory access was encountered
+// UNSUPPORTED: amdgcn-amd-amdhsa
+// UNSUPPORTED: nvptx64-nvidia-cuda
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: intelgpu
+
+#pragma omp requires unified_shared_memory
+
+#include <stdio.h>
+
+#pragma omp begin declare target
+double base[10] = {0};
+#pragma omp end declare target
+
+int main() {
+
+// close range covers base[1] and should properly
+// update the budder from the outer mapping.
+#pragma omp target map(close, alloc : base[1 : 9])
+ {
+ base[1] = 99.0;
+ }
+#pragma omp target update from(base)
+ // CHECK: base[1] = 99.000000 (expected 99.0)
+ printf("base[1] = %f (expected 99.0)\n", base[1]);
+ return 0;
+}
\ No newline at end of file
diff --git a/offload/test/unified_shared_memory/close_sub_section_global_oob.c b/offload/test/unified_shared_memory/close_sub_section_global_oob.c
new file mode 100644
index 0000000000000..eaf5ff8c74490
--- /dev/null
+++ b/offload/test/unified_shared_memory/close_sub_section_global_oob.c
@@ -0,0 +1,33 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+// REQUIRES: unified_shared_memory
+// UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
+
+// amdgpu runtime crash
+// Fails on nvptx with error: an illegal memory access was encountered
+// UNSUPPORTED: amdgcn-amd-amdhsa
+// UNSUPPORTED: nvptx64-nvidia-cuda
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: intelgpu
+
+#pragma omp requires unified_shared_memory
+
+#include <stdio.h>
+
+#pragma omp begin declare target
+double base[10] = {0};
+#pragma omp end declare target
+
+int main() {
+
+// close on sub-section base[2:8]; access to base[1] is outside the close
+// buffer
+#pragma omp target map(close, alloc : base[2 : 8])
+ {
+ base[1] = 99.0;
+ }
+#pragma omp target update from(base)
+ // CHECK: base[1] = 99.000000 (expected 99.0)
+ printf("base[1] = %f (expected 99.0)\n", base[1]);
+ return 0;
+}
\ No newline at end of file
diff --git a/offload/test/unified_shared_memory/declare_target_map_extent.c b/offload/test/unified_shared_memory/declare_target_map_extent.c
new file mode 100644
index 0000000000000..e44532d70274b
--- /dev/null
+++ b/offload/test/unified_shared_memory/declare_target_map_extent.c
@@ -0,0 +1,67 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+// REQUIRES: unified_shared_memory
+// UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
+
+// amdgpu runtime crash
+// Fails on nvptx with error: an illegal memory access was encountered
+// UNSUPPORTED: amdgcn-amd-amdhsa
+// UNSUPPORTED: nvptx64-nvidia-cuda
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: intelgpu
+
+// Mapping a declare-target aggregate, or a subsection of one that starts at its
+// beginning, must work under unified shared memory.
+//
+// If the runtime registers a mapping for the storage of a declare-target
+// variable, the extent it registers has to be the extent of the variable. Using
+// the size of the offload entry instead describes the device reference pointer,
+// i.e. sizeof(void *), and then a map of the real object looks like an attempt
+// to extend an existing, smaller mapping, which is rejected:
+//
+// explicit extension not allowed: host address specified is ... (80 bytes),
+// but device allocation maps to host at ... (8 bytes)
+//
+// That aborts the program rather than producing a wrong value, and it happens
+// for a plain map as well as for a close one.
+
+#include <stdio.h>
+
+#pragma omp requires unified_shared_memory
+
+#pragma omp begin declare target
+double base[10] = {0};
+#pragma omp end declare target
+
+int main() {
+ // A plain map of the whole declare-target array.
+#pragma omp target map(tofrom : base[0 : 10])
+ {
+ base[0] = 7.0;
+ }
+ // CHECK: plain whole-array map: base[0] = 7.000000
+ printf("plain whole-array map: base[0] = %f\n", base[0]);
+
+ // A close map of the whole array.
+#pragma omp target map(close, alloc : base[0 : 10])
+ {
+ base[3] = 5.0;
+ }
+#pragma omp target update from(base)
+ // CHECK: close whole-array map: base[3] = 5.000000
+ printf("close whole-array map: base[3] = %f\n", base[3]);
+
+ // A close map of a subsection that starts at the beginning of the array, so
+ // it overlaps whatever was registered for the variable itself.
+#pragma omp target map(close, alloc : base[0 : 9])
+ {
+ base[1] = 99.0;
+ }
+#pragma omp target update from(base)
+ // CHECK: close leading subsection: base[1] = 99.000000
+ printf("close leading subsection: base[1] = %f\n", base[1]);
+
+ // CHECK: Done!
+ printf("Done!\n");
+ return 0;
+}
diff --git a/offload/test/unified_shared_memory/declare_target_tracking.c b/offload/test/unified_shared_memory/declare_target_tracking.c
new file mode 100644
index 0000000000000..fc3d68b27a064
--- /dev/null
+++ b/offload/test/unified_shared_memory/declare_target_tracking.c
@@ -0,0 +1,65 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+// REQUIRES: unified_shared_memory
+// UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
+
+// amdgpu runtime crash
+// Fails on nvptx with error: an illegal memory access was encountered
+// UNSUPPORTED: amdgcn-amd-amdhsa
+// UNSUPPORTED: nvptx64-nvidia-cuda
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: intelgpu
+
+// A `declare target` variable is in the device data environment for its whole
+// extent, so omp_target_is_present must report it as present for every byte of
+// it, not just for its first few bytes.
+//
+// Under unified shared memory such a variable is represented on the device by a
+// reference pointer to the host storage, and the offload entry for it describes
+// that *pointer*: it gives neither the variable's address nor its extent.
+//
+// Under unified shared memory such a variable is represented on the device by a
+// reference pointer to the host storage, and code generation communicates the
+// variable's own extent so the runtime can register its storage.
+
+#include <omp.h>
+#include <stdio.h>
+
+#pragma omp requires unified_shared_memory
+
+#pragma omp begin declare target
+int scalar = 111;
+int arr[64] = {0};
+#pragma omp end declare target
+
+static int present(void *P) {
+ return omp_target_is_present(P, omp_get_default_device());
+}
+
+int main() {
+ // Make sure the device image (and with it the declare-target registration)
+ // has been loaded before querying presence.
+#pragma omp target
+ {
+ }
+
+ // CHECK: scalar present: 1
+ printf("scalar present: %d\n", present(&scalar));
+
+ // Every element of a declare-target array is present, including the last one.
+ // CHECK: arr present first/mid/last: 1 1 1
+ printf("arr present first/mid/last: %d %d %d\n", present(&arr[0]),
+ present(&arr[32]), present(&arr[63]));
+
+ int fails = 0;
+ for (int i = 0; i < 64; ++i)
+ if (!present(&arr[i]))
+ fails++;
+ // CHECK: arr present for all elements: Succeeded
+ printf("arr present for all elements: %s\n",
+ (fails == 0) ? "Succeeded" : "Failed");
+
+ // CHECK: Done!
+ printf("Done!\n");
+ return 0;
+}
>From c5bf90f1b403a14713f616d9043653202e2505eb Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Tue, 4 Aug 2026 22:21:43 -0700
Subject: [PATCH 9/9] [OFFLOAD] Give a pointer device storage before attaching
it, under USM
Pointer attachment assigns the corresponding pointer. Under unified shared memory
the corresponding storage may be the original storage, and then the assignment is
observable through the original pointer: the host can no longer use it, and once
the pointee's storage is released it is left dangling. A close mapping of a
pointee makes this reachable, since the pointee gets a device buffer while its
pointer may still be on the host path.
Give the entry holding the pointer a device allocation of its own instead, so
that the pointer has a corresponding pointer distinct from the original. The
whole entry is given storage rather than just the pointer, because the device
address of anything inside it is defined as the entry's device address plus an
offset -- for a structure member, or for a pointer inside an untyped byte range,
there is nothing else it could be derived from. Its contents are copied so that
whatever else it holds is present on the device.
The entry keeps its identity, so its reference count continues to govern its
lifetime and the existing shadow pointer still protects the original pointer from
a later from-transfer of the surrounding storage.
Giving an entry storage moves it, so anything already attached to a pointee
within it designates its previous device address. Those pointers are attached
again, which may require giving a further entry storage, so this repeats until
nothing is left; it terminates because an entry can only be given an allocation
once. Finding them needs the pointee-to-pointer direction, which the shadow
pointers do not provide, so attachment now records it in
MappingInfoTy::AttachedPointers. Only pointers that are actually attached are
recorded, so this stays empty for mappings that never involve attachment.
The optimization of leaving mappings on the host path is preserved: storage is
given to an entry only when a pointer inside it is about to be attached.
Adds tests for a pointer that is a structure member, two pointers to one pointee
of which only one has device storage, a host read while a pointer is attached, a
structure mapped as an untyped byte range, and a chain where one pointer lives
inside another's pointee.
Co-Authored-By: Claude Opus 5 (1M context) <noreply at anthropic.com>
---
offload/include/OpenMP/Mapping.h | 95 +++++++++++++++-
offload/libomptarget/OpenMP/Mapping.cpp | 74 +++++++++++-
offload/libomptarget/omptarget.cpp | 105 +++++++++++++++---
.../close_attach_host_deref.c | 80 +++++++++++++
.../close_attach_pointer_chain.c | 91 +++++++++++++++
.../close_attach_struct_byte_view.c | 73 ++++++++++++
.../close_ptee_attached_host_read.c | 93 ++++++++++++++++
.../close_ptee_struct_member.c | 85 ++++++++++++++
.../close_ptee_two_pointers.c | 83 ++++++++++++++
.../close_ptr_ptee_nested.c | 14 ++-
.../close_ptr_ptee_samedir.c | 12 +-
11 files changed, 777 insertions(+), 28 deletions(-)
create mode 100644 offload/test/unified_shared_memory/close_attach_host_deref.c
create mode 100644 offload/test/unified_shared_memory/close_attach_pointer_chain.c
create mode 100644 offload/test/unified_shared_memory/close_attach_struct_byte_view.c
create mode 100644 offload/test/unified_shared_memory/close_ptee_attached_host_read.c
create mode 100644 offload/test/unified_shared_memory/close_ptee_struct_member.c
create mode 100644 offload/test/unified_shared_memory/close_ptee_two_pointers.c
diff --git a/offload/include/OpenMP/Mapping.h b/offload/include/OpenMP/Mapping.h
index 4f20cc5490143..af9e6f3acd3d2 100644
--- a/offload/include/OpenMP/Mapping.h
+++ b/offload/include/OpenMP/Mapping.h
@@ -117,8 +117,31 @@ struct HostDataToTargetTy {
const uintptr_t HstPtrEnd; // non-inclusive.
const map_var_info_t HstPtrName; // Optional source name of mapped variable.
- const uintptr_t TgtAllocBegin; // allocated target memory
- const uintptr_t TgtPtrBegin; // mapped target memory = TgtAllocBegin + padding
+ // Not const: an entry whose storage is the host storage can later be given a
+ // device allocation, see giveDeviceAllocation().
+ uintptr_t TgtAllocBegin; // allocated target memory
+ uintptr_t TgtPtrBegin; // mapped target memory = TgtAllocBegin + padding
+
+ /// Whether this entry's storage is the host storage itself, i.e. it owns no
+ /// device allocation. That is the case for the entries recorded on the
+ /// unified-shared-memory host path.
+ bool isHostBacked() const { return TgtPtrBegin == HstPtrBegin; }
+
+ /// Give a host-backed entry a device allocation, so that it stops sharing
+ /// storage with the original. Used when a pointer inside this entry's storage
+ /// is about to be attached: attachment assigns the corresponding pointer, and
+ /// while the corresponding storage is the original storage that assignment
+ /// would be observable through the original pointer.
+ ///
+ /// The entry keeps its identity, so its reference count continues to govern
+ /// its lifetime, and any shadow pointer recorded for it is transferred back
+ /// to the host by the same data motion as for any other entry.
+ void giveDeviceAllocation(uintptr_t NewTgtAllocBegin,
+ uintptr_t NewTgtPtrBegin) {
+ assert(isHostBacked() && "Entry already owns a device allocation");
+ TgtAllocBegin = NewTgtAllocBegin;
+ TgtPtrBegin = NewTgtPtrBegin;
+ }
private:
static const uint64_t INFRefCount = ~(uint64_t)0;
@@ -495,6 +518,25 @@ struct AttachMapInfo {
MapType(Type), Pointername(Name) {}
};
+/// A pointer that has been attached to a pointee, recorded so that it can be
+/// attached again if the pointee's entry is later given a device allocation and
+/// its device address therefore changes.
+///
+/// The pointee is the key of MappingInfoTy::AttachedPointers, so only what is
+/// needed to redo the attachment is kept here.
+struct AttachedPointerTy {
+ /// Original address of the pointer itself, e.g. &p or &s.p.
+ void **HstPtrAddr;
+ /// Original address of the pointee base, which may differ from the pointee
+ /// address that keys the index, e.g. for map(p[10:5]).
+ void *HstPteeBase;
+ /// Size of the pointer, which is larger than a pointer for a descriptor.
+ int64_t PtrSize;
+
+ AttachedPointerTy(void **HstPtrAddr, void *HstPteeBase, int64_t PtrSize)
+ : HstPtrAddr(HstPtrAddr), HstPteeBase(HstPteeBase), PtrSize(PtrSize) {}
+};
+
/// Structure to track new allocations, ATTACH entries, DELETE entries and
/// skipped FROM data transfer information for a given construct, across
/// recursive calls (for handling mappers) to targetDataBegin/targetDataEnd.
@@ -665,6 +707,23 @@ struct MappingInfoTy {
/// The type used to access the HDTT map.
using HDTTMapAccessorTy = decltype(HostDataToTargetMap)::AccessorTy;
+ /// Which pointers are attached to a given pointee, keyed by the pointee's
+ /// original address.
+ ///
+ /// Attachment records a shadow pointer on the entry holding the pointer,
+ /// which gives the pointer's pointee but not the reverse. This index provides
+ /// the reverse direction, which is needed when a pointee's entry is given a
+ /// device allocation after something has already been attached to it: the
+ /// attached pointers designate its previous device address and have to be
+ /// attached again. See giveEntryDeviceAllocation().
+ ///
+ /// Only pointers that are actually attached are recorded, so this stays empty
+ /// for mappings that never involve pointer attachment.
+ ///
+ /// Accessed under the HDTT map accessor.
+ llvm::DenseMap<void *, llvm::SmallVector<AttachedPointerTy, 2>>
+ AttachedPointers;
+
/// Lookup the mapping of \p HstPtrBegin in \p HDTTMap. The accessor ensures
/// exclusive access to the HDTT map.
LookupResult lookupMapping(HDTTMapAccessorTy &HDTTMap, void *HstPtrBegin,
@@ -719,6 +778,38 @@ struct MappingInfoTy {
[[nodiscard]] int eraseMapEntry(HDTTMapAccessorTy &HDTTMap,
HostDataToTargetTy *Entry, int64_t Size);
+ /// Record that \p HstPtrAddr is attached to the pointee at \p HstPteeBegin.
+ /// See AttachedPointers.
+ void recordAttachedPointer(void *HstPteeBegin, void **HstPtrAddr,
+ void *HstPteeBase, int64_t PtrSize);
+
+ /// Give \p Entry, which must be host-backed, a device allocation and copy the
+ /// current contents of its storage into it, so that it stops sharing storage
+ /// with the original.
+ ///
+ /// This is needed before attaching a pointer that lies within \p Entry's
+ /// storage: attachment assigns the corresponding pointer, and while the
+ /// corresponding storage is the original storage that assignment is
+ /// observable through the original pointer. The whole entry has to be given
+ /// storage rather than just the pointer, because the device address of
+ /// anything inside it is defined as this entry's device address plus an
+ /// offset.
+ ///
+ /// The contents are copied so that whatever else the storage holds -- other
+ /// structure members, for instance -- is present on the device.
+ ///
+ /// Anything already attached to a pointee inside \p Entry designated its
+ /// previous device address, so those pointers have to be attached again. They
+ /// are appended to \p ToReattach as (pointee, pointer) pairs rather than
+ /// attached here, since performing an attachment is the caller's job.
+ ///
+ /// \p HDTTMap must be held by the caller. Returns \c OFFLOAD_SUCCESS if the
+ /// entry now owns a device allocation, and \c OFFLOAD_FAIL otherwise.
+ [[nodiscard]] int giveEntryDeviceAllocation(
+ HDTTMapAccessorTy &HDTTMap, HostDataToTargetTy *Entry,
+ AsyncInfoTy &AsyncInfo,
+ llvm::SmallVectorImpl<std::pair<void *, AttachedPointerTy>> &ToReattach);
+
/// Deallocate the \p Entry from the device memory and delete it. Return \c
/// OFFLOAD_SUCCESS if the deallocation operations executed successfully, and
/// return \c OFFLOAD_FAIL otherwise.
diff --git a/offload/libomptarget/OpenMP/Mapping.cpp b/offload/libomptarget/OpenMP/Mapping.cpp
index dbf88118000fa..ed0751a1027eb 100644
--- a/offload/libomptarget/OpenMP/Mapping.cpp
+++ b/offload/libomptarget/OpenMP/Mapping.cpp
@@ -561,6 +561,73 @@ int MappingInfoTy::eraseMapEntry(HDTTMapAccessorTy &HDTTMap,
return OFFLOAD_SUCCESS;
}
+void MappingInfoTy::recordAttachedPointer(void *HstPteeBegin, void **HstPtrAddr,
+ void *HstPteeBase, int64_t PtrSize) {
+ auto &Attached = AttachedPointers[HstPteeBegin];
+ for (const AttachedPointerTy &A : Attached)
+ if (A.HstPtrAddr == HstPtrAddr)
+ return;
+ Attached.emplace_back(HstPtrAddr, HstPteeBase, PtrSize);
+}
+
+int MappingInfoTy::giveEntryDeviceAllocation(
+ HDTTMapAccessorTy &HDTTMap, HostDataToTargetTy *Entry,
+ AsyncInfoTy &AsyncInfo,
+ llvm::SmallVectorImpl<std::pair<void *, AttachedPointerTy>> &ToReattach) {
+ assert(Entry && "Trying to allocate for a null entry.");
+ assert(Entry->isHostBacked() && "Entry already owns a device allocation");
+
+ void *HstPtrBegin = reinterpret_cast<void *>(Entry->HstPtrBegin);
+ int64_t Size = Entry->HstPtrEnd - Entry->HstPtrBegin;
+
+ uintptr_t TgtAllocBegin =
+ reinterpret_cast<uintptr_t>(Device.allocData(Size, HstPtrBegin));
+ if (!TgtAllocBegin) {
+ REPORT() << "Failed to allocate device memory for " << HstPtrBegin << ".";
+ return OFFLOAD_FAIL;
+ }
+
+ Entry->giveDeviceAllocation(TgtAllocBegin, TgtAllocBegin);
+
+ INFO(OMP_INFOTYPE_MAPPING_CHANGED, Device.DeviceID,
+ "Allocating device memory for existing map entry with "
+ "HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD ", Size=%" PRId64 "\n",
+ DPxPTR(Entry->HstPtrBegin), DPxPTR(TgtAllocBegin), Size);
+ ODBG(ODT_Mapping) << "Allocating device memory for existing map entry (hst:"
+ << HstPtrBegin
+ << ") -> (tgt:" << reinterpret_cast<void *>(TgtAllocBegin)
+ << "), Size=" << Size;
+
+ // The storage was shared with the host until now, so bring its current
+ // contents over: it may hold data other than the pointer being attached.
+ int Ret = Device.submitData(reinterpret_cast<void *>(TgtAllocBegin),
+ HstPtrBegin, Size, AsyncInfo, Entry);
+ if (Ret != OFFLOAD_SUCCESS) {
+ REPORT() << "Copying data to device failed.";
+ return OFFLOAD_FAIL;
+ }
+
+ if (Device.notifyDataMapped(HstPtrBegin, Size))
+ return OFFLOAD_FAIL;
+
+ // The entry's device address has changed, so anything already attached to a
+ // pointee within it designates the previous one. Collect those pointers for
+ // the caller to attach again.
+ for (auto &[HstPteeBegin, Attached] : AttachedPointers) {
+ if (reinterpret_cast<uintptr_t>(HstPteeBegin) < Entry->HstPtrBegin ||
+ reinterpret_cast<uintptr_t>(HstPteeBegin) >= Entry->HstPtrEnd)
+ continue;
+ for (const AttachedPointerTy &A : Attached) {
+ ODBG(ODT_Mapping) << "Pointer " << A.HstPtrAddr
+ << " is attached to pointee " << HstPteeBegin
+ << " within the entry, so it needs attaching again";
+ ToReattach.emplace_back(HstPteeBegin, A);
+ }
+ }
+
+ return OFFLOAD_SUCCESS;
+}
+
int MappingInfoTy::deallocTgtPtrAndEntry(HostDataToTargetTy *Entry,
int64_t Size) {
assert(Entry && "Trying to deallocate a null entry.");
@@ -577,11 +644,10 @@ int MappingInfoTy::deallocTgtPtrAndEntry(HostDataToTargetTy *Entry,
return OFFLOAD_FAIL;
}
- // The reuse entry recorded on the unified-shared-memory host path owns no
- // device allocation: its allocation maps to the host address itself, so it
- // must not be handed to deleteData().
+ // A host-backed entry owns no device allocation: its allocation maps to the
+ // host address itself, so it must not be handed to deleteData().
int Ret = OFFLOAD_SUCCESS;
- if (Entry->TgtAllocBegin != Entry->HstPtrBegin)
+ if (!Entry->isHostBacked())
Ret = Device.deleteData((void *)Entry->TgtAllocBegin);
// Notify the plugin about the unmapped memory.
diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp
index 917b865da53f5..88e12c96c8f77 100644
--- a/offload/libomptarget/omptarget.cpp
+++ b/offload/libomptarget/omptarget.cpp
@@ -453,6 +453,13 @@ static int performPointerAttachment(DeviceTy &Device, AsyncInfoTy &AsyncInfo,
void *TgtPteeBase =
calculateTargetPointeeBase(HstPteeBase, HstPteeBegin, TgtPteeBegin);
+ // Record the pointer against its pointee, so that it can be attached again if
+ // the pointee's entry is later given a device allocation. Recorded even when
+ // the shadow pointer below turns out to be a duplicate, since the pointee's
+ // address may still change afterwards.
+ Device.getMappingInfo().recordAttachedPointer(HstPteeBegin, HstPtrAddr,
+ HstPteeBase, HstPtrSize);
+
// Add shadow pointer tracking
if (!PtrTPR.getEntry()->addShadowPointer(
ShadowPtrInfoTy{HstPtrAddr, TgtPtrAddr, TgtPteeBase, HstPtrSize})) {
@@ -878,15 +885,9 @@ int processAttachEntries(DeviceTy &Device, StateInfoTy &StateInfo,
}
// Lambda to perform target pointer lookup and validation
- // \p AllowHostPointer permits the lookup to succeed for storage shared with
- // the host. That is correct for the pointer being attached: the device
- // dereferences the same storage, so attaching writes the device pointee
- // address into it and the original value is restored at the end of the
- // region through the shadow-pointer mechanism. It is not correct for the
- // pointee, where a host address would mean there is nothing to attach to.
auto LookupTargetPointer =
- [&](void *Ptr, int64_t Size, const char *PtrType,
- bool AllowHostPointer) -> std::optional<TargetPointerResultTy> {
+ [&](void *Ptr, int64_t Size,
+ const char *PtrType) -> std::optional<TargetPointerResultTy> {
// ATTACH map-type does not change ref-count, or do any allocation
// We just need to do a lookup for the pointer/pointee.
TargetPointerResultTy TPR = Device.getMappingInfo().getTgtPtrBegin(
@@ -902,7 +903,7 @@ int processAttachEntries(DeviceTy &Device, StateInfoTy &StateInfo,
<< PtrType << " not present on device";
return std::nullopt;
}
- if (TPR.Flags.IsHostPointer && !AllowHostPointer) {
+ if (TPR.Flags.IsHostPointer) {
ODBG(ODT_Mapping) << "Skipping ATTACH entry " << EntryIdx
<< ": device version of the " << PtrType
<< " is a host pointer.";
@@ -915,8 +916,7 @@ int processAttachEntries(DeviceTy &Device, StateInfoTy &StateInfo,
// Get device version of the pointee (e.g., &p[10]) first, as we can
// release its TPR after extracting the pointer value.
void *TgtPteeBegin = [&]() -> void * {
- if (auto PteeTPROpt = LookupTargetPointer(HstPteeBegin, 0, "pointee",
- /*AllowHostPointer=*/false))
+ if (auto PteeTPROpt = LookupTargetPointer(HstPteeBegin, 0, "pointee"))
return PteeTPROpt->TargetPointer;
return nullptr;
}();
@@ -926,10 +926,86 @@ int processAttachEntries(DeviceTy &Device, StateInfoTy &StateInfo,
// Get device version of the pointer (e.g., &p) next. We need to keep its
// TPR for use in shadow-pointer handling during pointer-attachment.
- auto PtrTPROpt = LookupTargetPointer(HstPtr, PtrSize, "pointer",
- /*AllowHostPointer=*/true);
+ auto PtrTPROpt = LookupTargetPointer(HstPtr, PtrSize, "pointer");
if (!PtrTPROpt)
continue;
+
+ // Attachment assigns the corresponding pointer. If the entry holding the
+ // pointer is host-backed then the corresponding storage *is* the original
+ // storage, so that assignment is observable through the original pointer:
+ // the host can no longer use it, and once the pointee's storage is released
+ // it is left dangling.
+ //
+ // Give that entry a device allocation of its own instead, so that the
+ // pointer has a corresponding pointer distinct from the original. The whole
+ // entry is given storage, since the device address of anything inside it is
+ // defined as the entry's device address plus an offset, and its contents
+ // are copied so that whatever else it holds is present on the device.
+ if (PtrTPROpt->getEntry() && PtrTPROpt->getEntry()->isHostBacked()) {
+ HostDataToTargetTy *Entry = PtrTPROpt->getEntry();
+ // Release the TPR's hold on the entry before taking the map accessor.
+ PtrTPROpt.reset();
+
+ // Giving an entry a device allocation moves it, so anything already
+ // attached to a pointee within it has to be attached again. Doing so may
+ // require giving a further entry an allocation, so keep going until
+ // nothing is left, which terminates because an entry can only be given an
+ // allocation once.
+ SmallVector<std::pair<void *, AttachedPointerTy>> ToReattach;
+ {
+ MappingInfoTy::HDTTMapAccessorTy HDTTMap =
+ Device.getMappingInfo().HostDataToTargetMap.getExclusiveAccessor();
+ if (Device.getMappingInfo().giveEntryDeviceAllocation(
+ HDTTMap, Entry, AsyncInfo, ToReattach) != OFFLOAD_SUCCESS)
+ return OFFLOAD_FAIL;
+ }
+
+ while (!ToReattach.empty()) {
+ auto [ReHstPteeBegin, ReAttached] = ToReattach.pop_back_val();
+
+ auto RePteeTPROpt =
+ LookupTargetPointer(ReHstPteeBegin, 0, "re-attach pointee");
+ if (!RePteeTPROpt)
+ continue;
+ void *ReTgtPteeBegin = RePteeTPROpt->TargetPointer;
+ RePteeTPROpt.reset();
+
+ auto RePtrTPROpt = LookupTargetPointer(
+ ReAttached.HstPtrAddr, ReAttached.PtrSize, "re-attach pointer");
+ if (!RePtrTPROpt)
+ continue;
+
+ if (RePtrTPROpt->getEntry() &&
+ RePtrTPROpt->getEntry()->isHostBacked()) {
+ HostDataToTargetTy *ReEntry = RePtrTPROpt->getEntry();
+ RePtrTPROpt.reset();
+ {
+ MappingInfoTy::HDTTMapAccessorTy HDTTMap =
+ Device.getMappingInfo()
+ .HostDataToTargetMap.getExclusiveAccessor();
+ if (Device.getMappingInfo().giveEntryDeviceAllocation(
+ HDTTMap, ReEntry, AsyncInfo, ToReattach) != OFFLOAD_SUCCESS)
+ return OFFLOAD_FAIL;
+ }
+ RePtrTPROpt = LookupTargetPointer(
+ ReAttached.HstPtrAddr, ReAttached.PtrSize, "re-attach pointer");
+ if (!RePtrTPROpt)
+ continue;
+ }
+
+ Ret = performPointerAttachment(
+ Device, AsyncInfo, ReAttached.HstPtrAddr, ReAttached.HstPteeBase,
+ ReHstPteeBegin,
+ reinterpret_cast<void **>(RePtrTPROpt->TargetPointer),
+ ReTgtPteeBegin, ReAttached.PtrSize, *RePtrTPROpt);
+ if (Ret != OFFLOAD_SUCCESS)
+ return OFFLOAD_FAIL;
+ }
+
+ PtrTPROpt = LookupTargetPointer(HstPtr, PtrSize, "pointer");
+ if (!PtrTPROpt)
+ continue;
+ }
TargetPointerResultTy &PtrTPR = *PtrTPROpt;
void **TgtPtrBase = reinterpret_cast<void **>(PtrTPR.TargetPointer);
@@ -1002,8 +1078,7 @@ postProcessingTargetDataEnd(DeviceTy *Device,
// with a zero reference count and a later map(close, ...) of the same
// storage finds it and concludes the data is already on the device.
const bool IsHostBackedEntry =
- TPR.getEntry() != nullptr &&
- TPR.getEntry()->TgtAllocBegin == TPR.getEntry()->HstPtrBegin;
+ TPR.getEntry() != nullptr && TPR.getEntry()->isHostBacked();
bool DelEntry = !TPR.isHostPointer() || IsHostBackedEntry;
// If the last element from the mapper (for end transfer args comes in
diff --git a/offload/test/unified_shared_memory/close_attach_host_deref.c b/offload/test/unified_shared_memory/close_attach_host_deref.c
new file mode 100644
index 0000000000000..cc05226625ea4
--- /dev/null
+++ b/offload/test/unified_shared_memory/close_attach_host_deref.c
@@ -0,0 +1,80 @@
+// RUN: %libomptarget-compile-generic
+// RUN: %libomptarget-run-generic 2>&1 | %fcheck-generic
+//
+// RUN: %libomptarget-compile-generic -DVIA_ALWAYS=1
+// RUN: env LIBOMPTARGET_TREAT_ATTACH_AUTO_AS_ALWAYS=1 \
+// RUN: %libomptarget-run-generic 2>&1 | %fcheck-generic
+
+// REQUIRES: unified_shared_memory
+// UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
+
+// amdgpu runtime crash
+// Fails on nvptx with error: an illegal memory access was encountered
+// UNSUPPORTED: amdgcn-amd-amdhsa
+// UNSUPPORTED: nvptx64-nvidia-cuda
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: intelgpu
+
+// Dereferencing a pointer on the host while its corresponding pointer is in the
+// attached state, under unified shared memory.
+//
+// The pointer is mapped without close, so its corresponding storage is the
+// original storage. The pointee is mapped with close, so it gets a device
+// buffer of its own, and pointer attachment then writes that device address
+// into the corresponding pointer -- which is the original pointer. A host
+// dereference afterwards goes through the device address.
+//
+// Two orderings are covered. In the default configuration the pointee is mapped
+// with close after the pointer, so the attachment is triggered by the pointee
+// becoming newly mapped. With VIA_ALWAYS the pointee already has device storage
+// and the attachment is triggered separately, by a map of the zero-length array
+// section under LIBOMPTARGET_TREAT_ATTACH_AUTO_AS_ALWAYS (OpenMP 6.0 has no
+// attach map-type-modifier for C/C++, so the environment variable stands in for
+// attach(always)).
+//
+// FIXME: the value checked below is the one produced today; the expected value
+// is given alongside it. Reading x[0] through p is the natural thing for a
+// program to do here, and it does not modify the pointer, so nothing in the
+// OpenMP 6.0 map clause restrictions forbids it.
+
+#include <stdio.h>
+
+#pragma omp requires unified_shared_memory
+
+int x[10];
+int *p = &x[0];
+
+int main() {
+ for (int i = 0; i < 10; ++i)
+ x[i] = 42;
+
+ // The unified_shared_memory requirement is registered when a device image is
+ // loaded, so the program needs a target region for it to take effect.
+#pragma omp target
+ {
+ }
+
+ // CHECK: before: p == &x[0]
+ printf("before: p %s &x[0]\n", p == &x[0] ? "==" : "!=");
+
+#if VIA_ALWAYS
+ // Pointer and pointee both get their storage first, then attachment is
+ // triggered on its own.
+#pragma omp target enter data map(alloc : p)
+#pragma omp target enter data map(close, alloc : x[0 : 10])
+#pragma omp target enter data map(alloc : p[0 : 0])
+#else
+ // The pointee is newly mapped with close, which triggers the attachment.
+#pragma omp target enter data map(alloc : p)
+#pragma omp target enter data map(close, alloc : p[0 : 10])
+#endif
+
+ // The pointer's entry was given a device allocation, so attachment wrote the
+ // device pointee address there and the original p is untouched.
+ // CHECK: after attach: p == &x[0]
+ printf("after attach: p %s &x[0]\n", p == &x[0] ? "==" : "!=");
+
+ // CHECK: Done!
+ printf("Done!\n");
+ return 0;
+}
diff --git a/offload/test/unified_shared_memory/close_attach_pointer_chain.c b/offload/test/unified_shared_memory/close_attach_pointer_chain.c
new file mode 100644
index 0000000000000..9678d8dba2233
--- /dev/null
+++ b/offload/test/unified_shared_memory/close_attach_pointer_chain.c
@@ -0,0 +1,91 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+// REQUIRES: unified_shared_memory
+// UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
+
+// amdgpu runtime crash
+// Fails on nvptx with error: an illegal memory access was encountered
+// UNSUPPORTED: amdgcn-amd-amdhsa
+// UNSUPPORTED: nvptx64-nvidia-cuda
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: intelgpu
+
+// A chain of two pointers, where the second one lives inside what the first one
+// points to: p1 refers to a structure whose member p2 refers to an array.
+//
+// p1 is attached to the structure first. Giving the structure device storage
+// afterwards -- because a close mapping of what p2 refers to needs p2 to have a
+// corresponding pointer distinct from the original -- moves the structure's
+// device address, so the value already attached to p1 no longer designates it.
+// Nothing revisits p1, so the device copy of p1 keeps pointing at the original
+// structure, and device code reaching p2 through p1 does not see the close
+// buffer.
+//
+// Any mechanism that gives an entry device storage after a pointer has been
+// attached to it therefore has to revisit the pointers already attached to that
+// entry, transitively.
+//
+// FIXME: the values checked below are the ones produced today; the expected
+// value is given alongside each.
+
+#include <stdio.h>
+
+#pragma omp requires unified_shared_memory
+
+struct Inner {
+ int *p2;
+ int pad;
+};
+
+int leaf[10];
+struct Inner inner;
+struct Inner *p1 = &inner;
+
+struct Inner *p1_device;
+int *p2_device;
+
+int main() {
+ for (int i = 0; i < 10; ++i)
+ leaf[i] = 42;
+ inner.p2 = &leaf[0];
+
+ // The unified_shared_memory requirement is registered when a device image is
+ // loaded, so the program needs a target region for it to take effect.
+#pragma omp target
+ {
+ }
+
+ // p1 is attached to the structure.
+#pragma omp target enter data map(alloc : p1)
+#pragma omp target enter data map(alloc : p1[0 : 1])
+
+ // The close mapping needs p2, which lives inside the structure, to have a
+ // corresponding pointer distinct from the original.
+#pragma omp target enter data map(close, alloc : p1->p2[0 : 10])
+
+#pragma omp target map(present, alloc : p1, p1[0 : 1]) \
+ map(from : p1_device, p2_device)
+ {
+ p1_device = p1;
+ p2_device = p1->p2;
+ }
+
+ // p1 was attached to the structure before it was given device storage, so it
+ // was attached again afterwards and designates the device copy.
+ // CHECK: device p1 != &inner
+ printf("device p1 %s &inner\n", p1_device == &inner ? "==" : "!=");
+
+ // So p2 is read from the device copy of the structure, where it designates
+ // the close buffer.
+ // CHECK: device p1->p2 != &leaf[0]
+ printf("device p1->p2 %s &leaf[0]\n", p2_device == &leaf[0] ? "==" : "!=");
+
+ // The original pointers are intact.
+ // CHECK: host: p1 == &inner, inner.p2 == &leaf[0]
+ printf("host: p1 %s &inner, inner.p2 %s &leaf[0]\n",
+ p1 == &inner ? "==" : "!=", inner.p2 == &leaf[0] ? "==" : "!=");
+
+ // CHECK: Done!
+ printf("Done!\n");
+ return 0;
+}
diff --git a/offload/test/unified_shared_memory/close_attach_struct_byte_view.c b/offload/test/unified_shared_memory/close_attach_struct_byte_view.c
new file mode 100644
index 0000000000000..907c4eb02c880
--- /dev/null
+++ b/offload/test/unified_shared_memory/close_attach_struct_byte_view.c
@@ -0,0 +1,73 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+// REQUIRES: unified_shared_memory
+// UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
+
+// amdgpu runtime crash
+// Fails on nvptx with error: an illegal memory access was encountered
+// UNSUPPORTED: amdgcn-amd-amdhsa
+// UNSUPPORTED: nvptx64-nvidia-cuda
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: intelgpu
+
+// A structure containing a pointer, mapped as an untyped byte range rather than
+// by its own name, and then a close mapping of what the member points to.
+//
+// The byte range is mapped without close, so its corresponding storage is the
+// original storage. Attachment for the member then writes the device pointee
+// address into the original member, since a member's device address is the
+// structure's device address plus the member offset.
+//
+// This is the same situation as close_ptee_struct_member.c, but the mapped
+// entry carries no type information: the runtime sees only a range of bytes
+// that happens to contain a pointer. Anything that gives the pointer storage of
+// its own would have to promote this whole range, which is the only storage the
+// member's device address can be derived from.
+//
+// FIXME: the value checked below is the one produced today; the expected value
+// is given alongside it.
+
+#include <stdio.h>
+
+#pragma omp requires unified_shared_memory
+
+int arr[10];
+
+struct S {
+ int x;
+ int y;
+ int *p;
+};
+
+struct S s;
+char *buf = (char *)&s;
+
+int main() {
+ for (int i = 0; i < 10; ++i)
+ arr[i] = 42;
+ s.p = &arr[0];
+
+ // The unified_shared_memory requirement is registered when a device image is
+ // loaded, so the program needs a target region for it to take effect.
+#pragma omp target
+ {
+ }
+
+ // CHECK: before: s.p == &arr[0]
+ printf("before: s.p %s &arr[0]\n", s.p == &arr[0] ? "==" : "!=");
+
+ // The structure is mapped as plain bytes, so it stays on the host path.
+#pragma omp target enter data map(alloc : buf[0 : sizeof(struct S)])
+
+ // The pointee is newly mapped with close, which triggers the attachment.
+#pragma omp target enter data map(close, alloc : s.p[0 : 10])
+
+ // The byte range was given a device allocation as a whole, so attachment
+ // wrote into that and the original s.p is untouched.
+ // CHECK: after attach: s.p == &arr[0]
+ printf("after attach: s.p %s &arr[0]\n", s.p == &arr[0] ? "==" : "!=");
+
+ // CHECK: Done!
+ printf("Done!\n");
+ return 0;
+}
diff --git a/offload/test/unified_shared_memory/close_ptee_attached_host_read.c b/offload/test/unified_shared_memory/close_ptee_attached_host_read.c
new file mode 100644
index 0000000000000..9394c646decd6
--- /dev/null
+++ b/offload/test/unified_shared_memory/close_ptee_attached_host_read.c
@@ -0,0 +1,93 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+// REQUIRES: unified_shared_memory
+// UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
+
+// amdgpu runtime crash
+// Fails on nvptx with error: an illegal memory access was encountered
+// UNSUPPORTED: amdgcn-amd-amdhsa
+// UNSUPPORTED: nvptx64-nvidia-cuda
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: intelgpu
+
+// The pointee already has device storage before the pointers are mapped, so no
+// allocation happens for it on the inner construct and there is nothing to
+// undo: the pointee is simply already on the device.
+//
+// p2 is mapped without close, so its corresponding storage is the original
+// storage and attachment writes the device pointee address into the original
+// p2. The host can then observe that value for as long as p2 remains attached,
+// which under OpenMP 6.0 lasts until the pointer's storage is removed from the
+// device data environment -- there is no detachment.
+//
+// This case is therefore not fixed by declining to allocate the pointee: it is
+// about attachment through shared storage, not about the close allocation.
+//
+// Note what OpenMP 6.0 does and does not say here. Attachment assigns the
+// corresponding pointer (7.9.6), and the corresponding storage may share
+// storage with the original (7.9.6, 1.3.2), in which case the assignment is
+// observable on the host. Nothing preserves the original value during that
+// window, and nothing restores it: the map-exiting sequence has no detach step.
+// The one place the specification confronts the same situation, for self maps,
+// requires runtime error termination when "the list item is a pointer that
+// would be assigned a different value as a result of pointer attachment"
+// (7.9.6), which suggests the intent is for this configuration not to arise
+// rather than for the host value to be preserved.
+//
+// FIXME: the values checked below are the ones produced today; the expected
+// value is given alongside each, on the reading that a program should be able
+// to use the original pointer while it happens to be attached.
+
+#include <stdio.h>
+
+#pragma omp requires unified_shared_memory
+
+int arr[10] = {0};
+
+int *p1 = &arr[0];
+int *p2 = &arr[0];
+
+int *p1_device, *p2_device;
+
+int main() {
+ // The pointee gets device storage before either pointer is mapped.
+#pragma omp target enter data map(close, alloc : arr[0 : 10])
+
+ // CHECK: before: p1 == &arr[0], p2 == &arr[0]
+ printf("before: p1 %s &arr[0], p2 %s &arr[0]\n",
+ p1 == &arr[0] ? "==" : "!=", p2 == &arr[0] ? "==" : "!=");
+
+#pragma omp target data map(close, alloc : p1, p1[0 : 0]) \
+ map(alloc : p2, p2[0 : 0])
+ {
+#pragma omp target map(present, alloc : p1, p2) map(from : p1_device, p2_device)
+ {
+ p1_device = p1;
+ p2_device = p2;
+ }
+
+ // Read on the host while both pointers are still attached.
+ //
+ // p1 has device storage of its own, so its original is unaffected.
+ // CHECK: inside: p1 == &arr[0]
+ printf("inside: p1 %s &arr[0]\n", p1 == &arr[0] ? "==" : "!=");
+
+ // CHECK: inside: p2 == &arr[0]
+ printf("inside: p2 %s &arr[0]\n", p2 == &arr[0] ? "==" : "!=");
+ }
+
+ // CHECK: in tgt: p1 != &arr[0], p2 != &arr[0]
+ printf(
+ "in tgt: p1 %s &arr[0], p2 %s &arr[0]\n",
+ p1_device == &arr[0] ? "==" : "!=", p2_device == &arr[0] ? "==" : "!=");
+
+ // CHECK: after: p1 == &arr[0]
+ printf("after: p1 %s &arr[0]\n", p1 == &arr[0] ? "==" : "!=");
+
+ // CHECK: after: p2 == &arr[0]
+ printf("after: p2 %s &arr[0]\n", p2 == &arr[0] ? "==" : "!=");
+
+ // CHECK: Done!
+ printf("Done!\n");
+ return 0;
+}
diff --git a/offload/test/unified_shared_memory/close_ptee_struct_member.c b/offload/test/unified_shared_memory/close_ptee_struct_member.c
new file mode 100644
index 0000000000000..aa512bc421a1e
--- /dev/null
+++ b/offload/test/unified_shared_memory/close_ptee_struct_member.c
@@ -0,0 +1,85 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+// REQUIRES: unified_shared_memory
+// UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
+
+// amdgpu runtime crash
+// Fails on nvptx with error: an illegal memory access was encountered
+// UNSUPPORTED: amdgcn-amd-amdhsa
+// UNSUPPORTED: nvptx64-nvidia-cuda
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: intelgpu
+
+// A close mapping of a pointee whose pointer is a structure member, where the
+// structure itself stays on the unified-shared-memory host path.
+//
+// Pointer attachment writes the device pointee address into the corresponding
+// pointer, which for a member is an interior address of the structure's own
+// storage: the device address of s.p is defined as the device address of s plus
+// the member offset. When s shares storage with the original, that write lands
+// in the original s.p, and device code that reaches the member through the
+// structure base -- as bar() does below -- observes it.
+//
+// The close pointee therefore gets a device buffer that the kernel writes,
+// while nothing copies it back to the original storage, so the write is lost.
+//
+// FIXME: the values checked below are the ones produced today; the expected
+// value is given alongside each. Note that the pointer cannot be given storage
+// of its own to attach into, precisely because a member's device address is
+// derived from the structure's: see the comment above.
+
+#include <stdio.h>
+
+#pragma omp requires unified_shared_memory
+
+int arr[10] = {0};
+
+struct S {
+ int a;
+ int *p;
+ int b;
+};
+
+struct S s = {1, &arr[0], 2};
+
+int *p_device;
+
+#pragma omp begin declare target
+// Reaches the member through the structure base, so it must see the attached
+// value at the device address of s plus the offset of p.
+void bar(struct S *ps) { ps->p[0] = 77; }
+#pragma omp end declare target
+
+int main() {
+ // CHECK: before: s.p == &arr[0]
+ printf("before: s.p %s &arr[0]\n", s.p == &arr[0] ? "==" : "!=");
+
+#pragma omp target data map(tofrom : s)
+ {
+#pragma omp target enter data map(close, to : s.p[0 : 10])
+
+#pragma omp target map(present, alloc : s) map(from : p_device)
+ {
+ p_device = s.p;
+ bar(&s);
+ }
+ }
+
+ // The member pointer was attached to the close buffer.
+ // CHECK: in tgt: s.p != &arr[0]
+ printf("in tgt: s.p %s &arr[0]\n", p_device == &arr[0] ? "==" : "!=");
+
+ // The original member pointer is intact afterwards.
+ // CHECK: after: s.p == &arr[0]
+ printf("after: s.p %s &arr[0]\n", s.p == &arr[0] ? "==" : "!=");
+
+ // bar() wrote through the attached member pointer, i.e. into the device
+ // buffer for the pointee. That buffer is mapped with alloc, which never
+ // assigns, so the value is not copied back.
+ // CHECK: arr[0] = 0
+ printf("arr[0] = %d\n", arr[0]);
+
+ // CHECK: Done!
+ printf("Done!\n");
+ return 0;
+}
diff --git a/offload/test/unified_shared_memory/close_ptee_two_pointers.c b/offload/test/unified_shared_memory/close_ptee_two_pointers.c
new file mode 100644
index 0000000000000..8b242292fcd59
--- /dev/null
+++ b/offload/test/unified_shared_memory/close_ptee_two_pointers.c
@@ -0,0 +1,83 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+// REQUIRES: unified_shared_memory
+// UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
+
+// amdgpu runtime crash
+// Fails on nvptx with error: an illegal memory access was encountered
+// UNSUPPORTED: amdgcn-amd-amdhsa
+// UNSUPPORTED: nvptx64-nvidia-cuda
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: intelgpu
+
+// Two pointers to the same pointee, where one of them has device storage of its
+// own and the other stays on the unified-shared-memory host path.
+//
+// p1 is mapped with close, so it gets a real device allocation and attachment
+// writes into that, leaving the original p1 untouched. p2 is mapped without
+// close, so its corresponding storage is the original storage, and attachment
+// writes the device pointee address into the original p2.
+//
+// So a single close pointee cannot be made correct by giving its pointers
+// storage of their own: whether that is possible depends on how each pointer
+// was mapped, and here one of them was not mapped in a way that provides it.
+//
+// FIXME: the values checked below are the ones produced today; the expected
+// value is given alongside each.
+
+#include <stdio.h>
+
+#pragma omp requires unified_shared_memory
+
+int arr[10] = {0};
+
+// Both point to the same pointee.
+int *p1 = &arr[0];
+int *p2 = &arr[0];
+
+int *p1_device, *p2_device;
+
+int main() {
+ // CHECK: before: p1 == &arr[0], p2 == &arr[0]
+ printf("before: p1 %s &arr[0], p2 %s &arr[0]\n",
+ p1 == &arr[0] ? "==" : "!=", p2 == &arr[0] ? "==" : "!=");
+
+ // p1 gets device storage of its own, p2 does not.
+#pragma omp target data map(close, alloc : p1) map(alloc : p2)
+ {
+ // The pointee is newly mapped here, with close, so it gets a device buffer.
+#pragma omp target data map(close, alloc : p1[0 : 10]) map(p2[0 : 0])
+ {
+#pragma omp target map(present, alloc : p1, p2) map(from : p1_device, p2_device)
+ {
+ p1_device = p1;
+ p2_device = p2;
+ p1[0] = 55;
+ }
+ }
+ }
+
+ // Both pointers were attached to the close buffer.
+ // CHECK: in tgt: p1 != &arr[0], p2 != &arr[0]
+ printf(
+ "in tgt: p1 %s &arr[0], p2 %s &arr[0]\n",
+ p1_device == &arr[0] ? "==" : "!=", p2_device == &arr[0] ? "==" : "!=");
+
+ // p1 has device storage of its own, so the original is intact.
+ // CHECK: after: p1 == &arr[0]
+ printf("after: p1 %s &arr[0]\n", p1 == &arr[0] ? "==" : "!=");
+
+ // p2 had no device storage of its own, so its entry was given one before the
+ // attachment, leaving the original intact as well.
+ // CHECK: after: p2 == &arr[0]
+ printf("after: p2 %s &arr[0]\n", p2 == &arr[0] ? "==" : "!=");
+
+ // The pointee is mapped with alloc, which never assigns, so the value written
+ // into the device buffer is not copied back.
+ // CHECK: arr[0] = 0
+ printf("arr[0] = %d\n", arr[0]);
+
+ // CHECK: Done!
+ printf("Done!\n");
+ return 0;
+}
diff --git a/offload/test/unified_shared_memory/close_ptr_ptee_nested.c b/offload/test/unified_shared_memory/close_ptr_ptee_nested.c
index 62d363efa8ed8..aff17d3bdb533 100644
--- a/offload/test/unified_shared_memory/close_ptr_ptee_nested.c
+++ b/offload/test/unified_shared_memory/close_ptr_ptee_nested.c
@@ -86,15 +86,21 @@ int main() {
// When the pointee has a device allocation of its own, the pointer is
// attached to it and its original host value is restored at the end of the
- // region. ATTACHED: ATTACH entry {{.*}} processed successfully ATTACHED:
- // Restoring host pointer
+ // region.
+ //
+ // clang-format off
+ // ATTACHED: ATTACH entry {{.*}} processed successfully
+ // ATTACHED: Restoring host pointer
+ // clang-format on
// ALL: Before tgt: p == p_host
// The pointer is attached to the close-allocated device pointee, so its
// device value differs from the host address.
+ // A pointer that gets attached is given device storage of its own first, so
+ // its address on the device always differs from the host address.
// V1: In tgt: p_device != p_host
- // V1: In tgt: paddr_device == &p_host
+ // V1: In tgt: paddr_device != &p_host
// p itself is close-allocated, giving the pointer variable its own device
// storage (so &p differs on the device); the pointee stays on the USM host
@@ -107,7 +113,7 @@ int main() {
// -- is attached to it, so its device value differs from the host address
// while the pointer variable itself stays on the host path (&p matches).
// V3: In tgt: p_device != p_host
- // V3: In tgt: paddr_device == &p_host
+ // V3: In tgt: paddr_device != &p_host
// The host pointer must be intact afterwards and the kernel's write must be
// visible on the host.
diff --git a/offload/test/unified_shared_memory/close_ptr_ptee_samedir.c b/offload/test/unified_shared_memory/close_ptr_ptee_samedir.c
index 78bdc326e1da3..5496062fd8c02 100644
--- a/offload/test/unified_shared_memory/close_ptr_ptee_samedir.c
+++ b/offload/test/unified_shared_memory/close_ptr_ptee_samedir.c
@@ -73,14 +73,20 @@ int main() {
// When the pointee has a device allocation of its own, the pointer is
// attached to it and its original host value is restored at the end of the
- // region. ATTACHED: ATTACH entry {{.*}} processed successfully ATTACHED:
- // Restoring host pointer
+ // region.
+ //
+ // clang-format off
+ // ATTACHED: ATTACH entry {{.*}} processed successfully
+ // ATTACHED: Restoring host pointer
+ // clang-format on
// ALL: Before tgt: p == p_host
// The close pointee is device-allocated and p (host path) is attached to it.
+ // A pointer that gets attached is given device storage of its own first, so
+ // its address on the device always differs from the host address.
// V1: In tgt: p_device != p_host
- // V1: In tgt: paddr_device == &p_host
+ // V1: In tgt: paddr_device != &p_host
// p is close-allocated (own device storage, so &p differs); the pointee is a
// plain USM mapping, so the device pointer value equals the host address.
More information about the llvm-commits
mailing list