[clang] [llvm] [WIP][OFFLOAD][OpenMP] Fix USM close mapping, and track declare-target storage under USM (PR #213214)
via cfe-commits
cfe-commits at lists.llvm.org
Fri Jul 31 00:12:26 PDT 2026
github-actions[bot] wrote:
<!--LLVM CODE FORMAT COMMENT: {clang-format}-->
:warning: C/C++ code formatter, clang-format found issues in your code. :warning:
<details>
<summary>
You can test this locally with the following command:
</summary>
``````````bash
git-clang-format --diff origin/main HEAD --extensions cpp,h,c -- offload/test/unified_shared_memory/check_tracking.c offload/test/unified_shared_memory/close.c offload/test/unified_shared_memory/close_alloc_declare_target_lifetime.c offload/test/unified_shared_memory/close_data_motion.c offload/test/unified_shared_memory/close_ptr_ptee_nested.c offload/test/unified_shared_memory/close_ptr_ptee_samedir.c offload/test/unified_shared_memory/close_sub_section.c offload/test/unified_shared_memory/close_sub_section_global.c offload/test/unified_shared_memory/close_sub_section_global_oob.c offload/test/unified_shared_memory/close_sub_section_oob.c offload/test/unified_shared_memory/declare_target_map_extent.c offload/test/unified_shared_memory/declare_target_tracking.c clang/test/OpenMP/declare_target_local_usm_codegen.cpp clang/test/OpenMP/nvptx_target_requires_unified_shared_memory.cpp llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp offload/include/OpenMP/Mapping.h offload/libomptarget/OpenMP/Mapping.cpp offload/libomptarget/PluginManager.cpp offload/libomptarget/omptarget.cpp --diff_from_common_commit
``````````
:warning:
The reproduction instructions above might return results for more than one PR
in a stack if you are using a stacked PR workflow. You can limit the results by
changing `origin/main` to the base branch/commit you want to compare against.
:warning:
</details>
<details>
<summary>
View the diff from clang-format here.
</summary>
``````````diff
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index e53e55a79..90ef9ea17 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -11998,8 +11998,7 @@ std::unique_ptr<CodeExtractor> DeviceSharedMemOutlineInfo::createCodeExtractor(
void OpenMPIRBuilder::createOffloadEntry(Constant *ID, Constant *Addr,
uint64_t Size, int32_t Flags,
GlobalValue::LinkageTypes,
- StringRef Name,
- uint64_t PointeeSize) {
+ StringRef Name, uint64_t PointeeSize) {
if (!Config.isGPU()) {
llvm::offloading::emitOffloadingEntry(
M, object::OffloadKind::OFK_OpenMP, ID,
@@ -12409,15 +12408,15 @@ 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
+ // 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.
+ // 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))
diff --git a/offload/include/OpenMP/Mapping.h b/offload/include/OpenMP/Mapping.h
index aeed70715..4f20cc549 100644
--- a/offload/include/OpenMP/Mapping.h
+++ b/offload/include/OpenMP/Mapping.h
@@ -562,8 +562,8 @@ public:
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 752d55104..dbf881180 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 09bccf743..3d3befe09 100644
--- a/offload/libomptarget/PluginManager.cpp
+++ b/offload/libomptarget/PluginManager.cpp
@@ -517,25 +517,24 @@ static int loadImagesOntoDevice(DeviceTy &Device) {
}
// 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
+ // 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.
+ // 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)";
+ 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,
diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp
index 48f7488ec..2de67fc22 100644
--- a/offload/libomptarget/omptarget.cpp
+++ b/offload/libomptarget/omptarget.cpp
@@ -886,9 +886,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> {
@@ -1003,9 +1003,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/close_alloc_declare_target_lifetime.c b/offload/test/unified_shared_memory/close_alloc_declare_target_lifetime.c
index ab1f063f5..80bd58d8a 100644
--- a/offload/test/unified_shared_memory/close_alloc_declare_target_lifetime.c
+++ b/offload/test/unified_shared_memory/close_alloc_declare_target_lifetime.c
@@ -12,7 +12,8 @@
// 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.
+// `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
@@ -20,11 +21,12 @@
// 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()
+// 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.
+// 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>
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 c135ef3a7..62d363efa 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 ae8cb7985..78bdc326e 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/declare_target_map_extent.c b/offload/test/unified_shared_memory/declare_target_map_extent.c
index 314e2fca4..e44532d70 100644
--- a/offload/test/unified_shared_memory/declare_target_map_extent.c
+++ b/offload/test/unified_shared_memory/declare_target_map_extent.c
@@ -16,14 +16,14 @@
// 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:
+// 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.
+// 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>
@@ -51,8 +51,8 @@ int main() {
// 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.
+ // 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;
diff --git a/offload/test/unified_shared_memory/declare_target_tracking.c b/offload/test/unified_shared_memory/declare_target_tracking.c
index 7bdbcbb24..fc3d68b27 100644
--- a/offload/test/unified_shared_memory/declare_target_tracking.c
+++ b/offload/test/unified_shared_memory/declare_target_tracking.c
@@ -3,7 +3,6 @@
// 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
``````````
</details>
https://github.com/llvm/llvm-project/pull/213214
More information about the cfe-commits
mailing list