[clang] [llvm] [WIP][OFFLOAD][OpenMP] Fix USM close mapping, and track declare-target storage under USM (PR #213214)
Abhinav Gaba via cfe-commits
cfe-commits at lists.llvm.org
Fri Jul 31 00:10:24 PDT 2026
https://github.com/abhinavgaba created https://github.com/llvm/llvm-project/pull/213214
Follow-on to #208122. This branch includes the changes from that PR, plus the
changes suggested for it, and will be rebased once it lands.
Three commits:
1. **Fix USM mapping tracking without disabling device data motion.** This is the
change suggested on #208122. 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, so device writes were dropped. Instead the reuse entry
recorded on the USM host path is reclaimed with the region that created it, and
only its device deallocation is skipped, since its allocation maps to the host
address. The close reuse is gated on the containment flags alone.
2. **Attach pointers that are newly mapped on the USM host path.** Attachment is
governed by a list item's mapping being new for the construct, not by device
memory having been allocated for it, so a pointer mapped for the first time
under USM was never attached. This matters when a `close` mapping gives a
pointee its own device allocation while its pointer is mapped without `close`.
3. **Communicate declare-target variable size under USM.** 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 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.
This last commit is why the first one can leave declare-target alone. Without the
variable's extent the runtime has neither its address nor its size, so it cannot
register the storage, and a mapping of the variable, or of part of it, is not
recognized as referring to storage that is already on the device. With the extent
available the existing containment check covers those cases, 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 reuses the existing device buffer instead of
getting its own, which fixes the lost write when the mapping covers part of the
object or is reached through a pointer into it.
- `close_sub_section_global.c` and `close_sub_section_global_oob.c` pass, the
latter of which aborts today with `double free or corruption`, on main as well as
on #208122.
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 before; 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. Note that the MLIR and Flang test suites have not
been run against this yet.
Testing: `offload/test/x86_64-unknown-linux-gnu` and `clang/test/OpenMP` are clean,
with only the two pre-existing XFAILs in the offload suite.
🤖 Generated with [Claude Code](https://claude.com/claude-code)
>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/8] 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/8] 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/8] 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/8] 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/8] 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/8] [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/8] [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 484aeaa13f3da49d5918b9dc745ec4e978f89597 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/8] [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 | 34 +++++++--
offload/libomptarget/PluginManager.cpp | 31 ++++++++
.../unified_shared_memory/check_tracking.c | 38 ++++++++--
.../close_alloc_declare_target_lifetime.c | 72 +++++++++++++++++++
.../close_sub_section_global.c | 33 +++++++++
.../close_sub_section_global_oob.c | 33 +++++++++
.../declare_target_map_extent.c | 67 +++++++++++++++++
.../declare_target_tracking.c | 66 +++++++++++++++++
11 files changed, 380 insertions(+), 17 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..5f68ee0f7a1d5 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -11874,11 +11874,13 @@ 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 +12053,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 +12227,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 +12284,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 +12747,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 +12762,7 @@ void OffloadEntriesInfoManager::registerDeviceGlobalVarEntryInfo(
return;
}
Entry.setVarSize(VarSize);
+ Entry.setPointeeSize(PointeeSize);
Entry.setLinkage(Linkage);
Entry.setAddress(Addr);
} else {
@@ -12764,6 +12785,7 @@ void OffloadEntriesInfoManager::registerDeviceGlobalVarEntryInfo(
else
OffloadEntriesDeviceGlobalVar.try_emplace(
VarName, OffloadingEntriesNum, Addr, VarSize, Flags, Linkage, "");
+ OffloadEntriesDeviceGlobalVar[VarName].setPointeeSize(PointeeSize);
++OffloadingEntriesNum;
}
}
diff --git a/offload/libomptarget/PluginManager.cpp b/offload/libomptarget/PluginManager.cpp
index 41b653a60adfd..09bccf7437dcc 100644
--- a/offload/libomptarget/PluginManager.cpp
+++ b/offload/libomptarget/PluginManager.cpp
@@ -516,6 +516,37 @@ 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/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..ab1f063f50f27
--- /dev/null
+++ b/offload/test/unified_shared_memory/close_alloc_declare_target_lifetime.c
@@ -0,0 +1,72 @@
+// 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_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..314e2fca4a0e2
--- /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..7bdbcbb24ee15
--- /dev/null
+++ b/offload/test/unified_shared_memory/declare_target_tracking.c
@@ -0,0 +1,66 @@
+// 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;
+}
More information about the cfe-commits
mailing list