[OpenMP] libomptarget: Don't map alignment padding to host

[jdenny]

Summary

In the case of partially mapped structs, libomptarget sometimes adds
padding to device allocations to ensure they are aligned properly.
However, without this patch, it considers that padding to be mapped to
the host, which can cause presence checks (e.g.,
omp_target_is_present or a present modifier) to misbehave for
unmapped parts of the struct. This patch keeps the padding but treats
it as unmapped. See the new test case for examples.

[jdenny]

Ping.

[grokos]

I understand what you are trying to achieve here and indeed this patch does exactly what you want to. But my question is, do we really want to allow this behavior? The test case which was added to demonstrate what problem this patch fixes looks illegal from an OpenMP perspective. It is a case of "explicitly mapping more" of the same object, which is forbidden in OpenMP. I wasn't sure whether that restriction has been relaxed and it slipped under my radar, but the latest TR11 spec still says about map clauses (https://www.openmp.org/wp-content/uploads/openmp-TR11.pdf, pg. 159, lines 29-32):

If a list item is an element of a structure, and a different element of the structure has a
corresponding list item in the device data environment prior to a task encountering the
construct associated with the map clause, then the list item must also have a corresponding
list item in the device data environment prior to the task encountering the construct.

In other words, if you want to map a struct member via a later map clause and other members of the same struct have been mapped via an earlier map clause, then the former member should also have been mapped via the earlier map clause.

If there is a valid scenario that doesn't work without this patch then:

1. The valid scenario should be added as a new test case to highlight what this patch fixes
2. Modify this patch so that the illegal scenario in padding_not_mapped.c still fails, as libomptarget correctly detects the attempt to "explicitly map more of the same struct"

[jdenny]

In D149685#4332051, @grokos wrote:

It is a case of "explicitly mapping more" of the same object, which is forbidden in OpenMP.

Indeed, I misunderstood. Sorry for the noise, and thanks for explaining. I'll abandon the patch.

If there is a valid scenario that doesn't work without this patch

I'm not sure there is.

Modify this patch so that the illegal scenario in padding_not_mapped.c still fails, as libomptarget correctly detects the attempt to "explicitly map more of the same struct"

For the record, while I designed that test to trigger detection, libomptarget does not always detect such cases. For example, the second pair of directives in the existing test openmp/libomptarget/test/mapping/low_alignment.c seem to be designed to avoid detection. In other cases, detection depends on the amount of padding on the device, which depends on how the host memory happens to be aligned, which can vary from run to run of the same program. That is, detection is non-deterministic.

Now that I have a better handle on what's happening, it seems to me that the second pair of directives in openmp/libomptarget/test/mapping/low_alignment.c should be removed. I think that test only passes when all of the following conditions are true, but OpenMP doesn't guarantee them to be true:

1. The runtime doesn't diagnose the restriction violation because map(from : t.i, t.j) has too little device padding for map(from : t.a) to collide with it.

2. t.i = 21; and t.j = 31; do not produce memory access violations even though they write to the offsets of t.i and t.j computed from the base device address of t from map(from : t.a) and even though those offsets are not guaranteed to be allocated on the device.

3. The desired values for t.i and t.j are copied back to the host even though t.i = 21; and t.j = 31; do *not* write to the offset of t.i and t.j computed from the base device address of t from map(from : t.i, t.j). The reason the correct values are copied back is that the previous kernel wrote the same values to s.i and s.j, which happened to have the same device addresses.

Agreed?

[jdenny]

In D149685#4332051, @grokos wrote:

If there is a valid scenario that doesn't work without this patch then:

After further thought, I believe there are several such scenarios. For a struct member that only maps to device padding and thus shouldn't be considered present, omp target update should have no effect, omp_target_is_present should return false, and present modifiers should produce runtime errors. Did I miss anything in the spec that makes these scenarios invalid?

I updated the test and patch summary for these scenarios, and I rebased.

[jdenny]

Ping

[jdoerfert]

I think the motivation is valid. The code looks good, though it's early for me. It seems to pass all our alignment tests, which is a good start. LG

[jdenny]

Thanks for the review.

[cseo]

I think there's a corner case you may have missed. I'm still seeing this on aarch64:

/llvm-project/openmp/libomptarget/test/mapping/target_derefence_array_pointrs.cpp:28:12: error: CHECK: expected string not found in input
 // CHECK: 3
           ^
<stdin>:1:103: note: scanning from here
Libomptarget message: explicit extension not allowed: host address specified is 0x0000ffffd26ff7c0 (12 bytes), but device allocation maps to host at 0x0000ffffd26ff7c0 (8 bytes)
                                                                                                      ^
<stdin>:1:104: note: possible intended match here
Libomptarget message: explicit extension not allowed: host address specified is 0x0000ffffd26ff7c0 (12 bytes), but device allocation maps to host at 0x0000ffffd26ff7c0 (8 bytes)

Your patch fixes most failures, except this one. Fails with and without LTO enabled. Could you please check?

[jdenny]

In D149685#4471532, @cseo wrote:

I think there's a corner case you may have missed. I'm still seeing this on aarch64:

/llvm-project/openmp/libomptarget/test/mapping/target_derefence_array_pointrs.cpp:28:12: error: CHECK: expected string not found in input
 // CHECK: 3
           ^
<stdin>:1:103: note: scanning from here
Libomptarget message: explicit extension not allowed: host address specified is 0x0000ffffd26ff7c0 (12 bytes), but device allocation maps to host at 0x0000ffffd26ff7c0 (8 bytes)
                                                                                                      ^
<stdin>:1:104: note: possible intended match here
Libomptarget message: explicit extension not allowed: host address specified is 0x0000ffffd26ff7c0 (12 bytes), but device allocation maps to host at 0x0000ffffd26ff7c0 (8 bytes)

Your patch fixes most failures, except this one. Fails with and without LTO enabled. Could you please check?

I'm not testing on aarch64, but that test fails for me both before and after this patch. This patch is about padding generated for a partial struct mapping, which is irrelevant to that test, as far as I can tell.

[cseo]

In D149685#4474512, @jdenny wrote:

In D149685#4471532, @cseo wrote:

I think there's a corner case you may have missed. I'm still seeing this on aarch64:

/llvm-project/openmp/libomptarget/test/mapping/target_derefence_array_pointrs.cpp:28:12: error: CHECK: expected string not found in input
 // CHECK: 3
           ^
<stdin>:1:103: note: scanning from here
Libomptarget message: explicit extension not allowed: host address specified is 0x0000ffffd26ff7c0 (12 bytes), but device allocation maps to host at 0x0000ffffd26ff7c0 (8 bytes)
                                                                                                      ^
<stdin>:1:104: note: possible intended match here
Libomptarget message: explicit extension not allowed: host address specified is 0x0000ffffd26ff7c0 (12 bytes), but device allocation maps to host at 0x0000ffffd26ff7c0 (8 bytes)

Your patch fixes most failures, except this one. Fails with and without LTO enabled. Could you please check?

I'm not testing on aarch64, but that test fails for me both before and after this patch. This patch is about padding generated for a partial struct mapping, which is irrelevant to that test, as far as I can tell.

OK, thanks! I'll take a look then.

[jdenny]

In D149685#4479109, @cseo wrote:

OK, thanks! I'll take a look then.

If you want to fix target_derefence_array_pointrs.cpp, I suggest looking at its function zoo. It explicitly maps the pointer sa when mapping *sa is needed. Its mapping has memory aliasing between sa and sa->sa. I'm not sure what the test intended, and I'm not confident how the implementation is meant to handle those cases, so you might want to contact the test's author. (It would also be nice to fix the typos in the test name while you're there.)

openmp/libomptarget/include/device.h

 
 /// Map between host data and target data.
 struct HostDataToTargetTy {
   const uintptr_t HstPtrBase; // host info.
   const uintptr_t HstPtrBegin;
   const uintptr_t HstPtrEnd;       // non-inclusive.
   const map_var_info_t HstPtrName; // Optional source name of mapped variable.
 
-  const uintptr_t TgtPtrBegin; // target info.
+  const uintptr_t TgtAllocBegin; // allocated target memory
+  const uintptr_t TgtPtrBegin; // mapped target memory = TgtAllocBegin + padding
 
 private:
   static const uint64_t INFRefCount = ~(uint64_t)0;
   static std::string refCountToStr(uint64_t RefCount) {
     return RefCount == INFRefCount ? "INF" : std::to_string(RefCount);
   }
 
   struct StatesTy {
struct StatesTy {
     /// references this entry will actually delete it.
     int32_t DataEndThreadCount = 0;
   };
   // When HostDataToTargetTy is used by std::set, std::set::iterator is const
   // use unique_ptr to make States mutable.
   const std::unique_ptr<StatesTy> States;
 
 public:
-  HostDataToTargetTy(uintptr_t BP, uintptr_t B, uintptr_t E, uintptr_t TB,
+  HostDataToTargetTy(uintptr_t BP, uintptr_t B, uintptr_t E,
+                     uintptr_t TgtAllocBegin, uintptr_t TgtPtrBegin,
                      bool UseHoldRefCount, map_var_info_t Name = nullptr,
                      bool IsINF = false)
       : HstPtrBase(BP), HstPtrBegin(B), HstPtrEnd(E), HstPtrName(Name),
-        TgtPtrBegin(TB), States(std::make_unique<StatesTy>(UseHoldRefCount ? 0
-                                                           : IsINF ? INFRefCount
-                                                                   : 1,
-                                                           !UseHoldRefCount ? 0
-                                                           : IsINF ? INFRefCount
-                                                                   : 1)) {}
+        TgtAllocBegin(TgtAllocBegin), TgtPtrBegin(TgtPtrBegin),
+        States(std::make_unique<StatesTy>(UseHoldRefCount ? 0
+                                          : IsINF         ? INFRefCount
+                                          : IsINF          ? INFRefCount
 
   /// Get the total reference count.  This is smarter than just getDynRefCount()
   /// + getHoldRefCount() because it handles the case where at least one is
   /// infinity and the other is non-zero.
   uint64_t getTotalRefCount() const {
     if (States->DynRefCount == INFRefCount ||
         States->HoldRefCount == INFRefCount)
       return INFRefCount;
struct DeviceTy {
   /// exist, and if unified shared memory is not enabled, a new mapping will be
   /// created and the data will also be transferred accordingly. nullptr will be
   /// returned because of any of following reasons:
   /// - Data allocation failed;
   /// - The user tried to do an illegal mapping;
   /// - Data transfer issue fails.
   TargetPointerResultTy getTargetPointer(
       HDTTMapAccessorTy &HDTTMap, void *HstPtrBegin, void *HstPtrBase,
-      int64_t Size, map_var_info_t HstPtrName, bool HasFlagTo,
-      bool HasFlagAlways, bool IsImplicit, bool UpdateRefCount,
+      int64_t TgtPadding, int64_t Size, map_var_info_t HstPtrName,
+      bool HasFlagTo, bool HasFlagAlways, bool IsImplicit, bool UpdateRefCount,
       bool HasCloseModifier, bool HasPresentModifier, bool HasHoldModifier,
       AsyncInfoTy &AsyncInfo, HostDataToTargetTy *OwnedTPR = nullptr,
       bool ReleaseHDTTMap = true);
 
   /// Return the target pointer for \p HstPtrBegin in \p HDTTMap. The accessor
   /// ensures exclusive access to the HDTT map.
   void *getTgtPtrBegin(HDTTMapAccessorTy &HDTTMap, void *HstPtrBegin,
                        int64_t Size);

openmp/libomptarget/src/device.cpp

int DeviceTy::associatePtr(void *HstPtrBegin, void *TgtPtrBegin, int64_t Size) {
 
   // Mapping does not exist, allocate it with refCount=INF
   const HostDataToTargetTy &NewEntry =
       *HDTTMap
            ->emplace(new HostDataToTargetTy(
                /*HstPtrBase=*/(uintptr_t)HstPtrBegin,
                /*HstPtrBegin=*/(uintptr_t)HstPtrBegin,
                /*HstPtrEnd=*/(uintptr_t)HstPtrBegin + Size,
+               /*TgtAllocBegin=*/(uintptr_t)TgtPtrBegin,
                /*TgtPtrBegin=*/(uintptr_t)TgtPtrBegin,
                /*UseHoldRefCount=*/false, /*Name=*/nullptr,
                /*IsRefCountINF=*/true))
            .first->HDTT;
   DP("Creating new map entry: HstBase=" DPxMOD ", HstBegin=" DPxMOD
      ", HstEnd=" DPxMOD ", TgtBegin=" DPxMOD ", DynRefCount=%s, "
      "HoldRefCount=%s\n",
      DPxPTR(NewEntry.HstPtrBase), DPxPTR(NewEntry.HstPtrBegin),
if (Size == 0) {
     }
   }
 
   return LR;
 }
 
 TargetPointerResultTy DeviceTy::getTargetPointer(
     HDTTMapAccessorTy &HDTTMap, void *HstPtrBegin, void *HstPtrBase,
-    int64_t Size, map_var_info_t HstPtrName, bool HasFlagTo, bool HasFlagAlways,
-    bool IsImplicit, bool UpdateRefCount, bool HasCloseModifier,
-    bool HasPresentModifier, bool HasHoldModifier, AsyncInfoTy &AsyncInfo,
-    HostDataToTargetTy *OwnedTPR, bool ReleaseHDTTMap) {
+    int64_t TgtPadding, int64_t Size, map_var_info_t HstPtrName, bool HasFlagTo,
+    bool HasFlagAlways, bool IsImplicit, bool UpdateRefCount,
+    bool HasCloseModifier, bool HasPresentModifier, bool HasHoldModifier,
+    AsyncInfoTy &AsyncInfo, HostDataToTargetTy *OwnedTPR, bool ReleaseHDTTMap) {
 
   LookupResult LR = lookupMapping(HDTTMap, HstPtrBegin, Size, OwnedTPR);
   LR.TPR.Flags.IsPresent = true;
 
   // Release the mapping table lock only after the entry is locked by
   // attaching it to TPR. Once TPR is destroyed it will release the lock
   // on entry. If it is returned the lock will move to the returned object.
   // If LR.Entry is already owned/locked we avoid trying to lock it again.
DP("Mapping required by 'present' map type modifier does not exist for "
        "HstPtrBegin=" DPxMOD ", Size=%" PRId64 "\n",
        DPxPTR(HstPtrBegin), Size);
     MESSAGE("device mapping required by 'present' map type modifier does not "
             "exist for host address " DPxMOD " (%" PRId64 " bytes)",
             DPxPTR(HstPtrBegin), Size);
   } else if (Size) {
     // If it is not contained and Size > 0, we should create a new entry for it.
     LR.TPR.Flags.IsNewEntry = true;
-    uintptr_t Ptr = (uintptr_t)allocData(Size, HstPtrBegin);
+    uintptr_t TgtAllocBegin =
+        (uintptr_t)allocData(TgtPadding + Size, HstPtrBegin);
+    uintptr_t TgtPtrBegin = TgtAllocBegin + TgtPadding;
     // Release the mapping table lock only after the entry is locked by
     // attaching it to TPR.
     LR.TPR.setEntry(HDTTMap
                         ->emplace(new HostDataToTargetTy(
                             (uintptr_t)HstPtrBase, (uintptr_t)HstPtrBegin,
-                            (uintptr_t)HstPtrBegin + Size, Ptr, HasHoldModifier,
-                            HstPtrName))
+                            (uintptr_t)HstPtrBegin + Size, TgtAllocBegin,
+                            TgtPtrBegin, HasHoldModifier, HstPtrName))
                         .first->HDTT);
     INFO(OMP_INFOTYPE_MAPPING_CHANGED, DeviceID,
          "Creating new map entry with HstPtrBase=" DPxMOD
-         ", HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD ", Size=%ld, "
-         "DynRefCount=%s, HoldRefCount=%s, Name=%s\n",
-         DPxPTR(HstPtrBase), DPxPTR(HstPtrBegin), DPxPTR(Ptr), Size,
+         ", HstPtrBegin=" DPxMOD ", TgtAllocBegin=" DPxMOD
+         ", TgtPtrBegin=" DPxMOD
+         ", Size=%ld, DynRefCount=%s, HoldRefCount=%s, Name=%s\n",
+         DPxPTR(HstPtrBase), DPxPTR(HstPtrBegin), DPxPTR(TgtAllocBegin),
+         DPxPTR(TgtPtrBegin), Size,
          LR.TPR.getEntry()->dynRefCountToStr().c_str(),
          LR.TPR.getEntry()->holdRefCountToStr().c_str(),
          (HstPtrName) ? getNameFromMapping(HstPtrName).c_str() : "unknown");
-    LR.TPR.TargetPointer = (void *)Ptr;
+    LR.TPR.TargetPointer = (void *)TgtPtrBegin;
 
     // Notify the plugin about the new mapping.
     if (notifyDataMapped(HstPtrBegin, Size))
       return {{false /* IsNewEntry */, false /* IsHostPointer */},
               nullptr /* Entry */,
               nullptr /* TargetPointer */};
   } else {
     // This entry is not present and we did not create a new entry for it.
int DeviceTy::eraseMapEntry(HDTTMapAccessorTy &HDTTMap,
   }
 
   return OFFLOAD_SUCCESS;
 }
 
 int DeviceTy::deallocTgtPtrAndEntry(HostDataToTargetTy *Entry, int64_t Size) {
   assert(Entry && "Trying to deallocate a null entry.");
 
-  DP("Deleting tgt data " DPxMOD " of size %" PRId64 "\n",
-     DPxPTR(Entry->TgtPtrBegin), Size);
+  DP("Deleting tgt data " DPxMOD " of size %" PRId64 " by freeing allocation "
+     "starting at " DPxMOD "\n",
+     DPxPTR(Entry->TgtPtrBegin), Size, DPxPTR(Entry->TgtAllocBegin));
 
   void *Event = Entry->getEvent();
   if (Event && destroyEvent(Event) != OFFLOAD_SUCCESS) {
     REPORT("Failed to destroy event " DPxMOD "\n", DPxPTR(Event));
     return OFFLOAD_FAIL;
   }
 
-  int Ret = deleteData((void *)Entry->TgtPtrBegin);
+  int Ret = deleteData((void *)Entry->TgtAllocBegin);
 
   // Notify the plugin about the unmapped memory.
   Ret |= notifyDataUnmapped((void *)Entry->HstPtrBegin);
 
   delete Entry;
 
   return Ret;
 }
__tgt_target_table *DeviceTy::loadBinary(void *Img) {
   std::lock_guard<decltype(RTL->Mtx)> LG(RTL->Mtx);
   return RTL->load_binary(RTLDeviceID, Img);
 }
 
 void *DeviceTy::allocData(int64_t Size, void *HstPtr, int32_t Kind) {
   return RTL->data_alloc(RTLDeviceID, Size, HstPtr, Kind);
 }
 
-int32_t DeviceTy::deleteData(void *TgtPtrBegin, int32_t Kind) {
-  return RTL->data_delete(RTLDeviceID, TgtPtrBegin, Kind);
+int32_t DeviceTy::deleteData(void *TgtAllocBegin, int32_t Kind) {
+  return RTL->data_delete(RTLDeviceID, TgtAllocBegin, Kind);
 }
 
 static void printCopyInfo(int DeviceId, bool H2D, void *SrcPtrBegin,
                           void *DstPtrBegin, int64_t Size,
                           HostDataToTargetTy *HT) {
 
   INFO(OMP_INFOTYPE_DATA_TRANSFER, DeviceId,
        "Copying data from %s to %s, %sPtr=" DPxMOD ", %sPtr=" DPxMOD

openmp/libomptarget/src/omptarget.cpp

for (auto *HostEntriesBegin : PM->HostEntriesBeginRegistrationOrder) {
              "\n",
              DPxPTR(CurrHostEntry->addr), DPxPTR(CurrDeviceEntry->addr),
              CurrDeviceEntry->size);
           HDTTMap->emplace(new HostDataToTargetTy(
               (uintptr_t)CurrHostEntry->addr /*HstPtrBase*/,
               (uintptr_t)CurrHostEntry->addr /*HstPtrBegin*/,
               (uintptr_t)CurrHostEntry->addr +
                   CurrHostEntry->size /*HstPtrEnd*/,
+              (uintptr_t)CurrDeviceEntry->addr /*TgtAllocBegin*/,
               (uintptr_t)CurrDeviceEntry->addr /*TgtPtrBegin*/,
               false /*UseHoldRefCount*/, CurrHostEntry->name,
               true /*IsRefCountINF*/));
 
           // Notify about the new mapping.
           if (Device.notifyDataMapped(CurrHostEntry->addr, CurrHostEntry->size))
             return OFFLOAD_FAIL;
         }
for (int32_t I = 0; I < ArgNum; ++I) {
     void *HstPtrBegin = Args[I];
     void *HstPtrBase = ArgsBase[I];
     int64_t DataSize = ArgSizes[I];
     map_var_info_t HstPtrName = (!ArgNames) ? nullptr : ArgNames[I];
 
     // Adjust for proper alignment if this is a combined entry (for structs).
     // Look at the next argument - if that is MEMBER_OF this one, then this one
     // is a combined entry.
-    int64_t Padding = 0;
+    int64_t TgtPadding = 0;
     const int NextI = I + 1;
     if (getParentIndex(ArgTypes[I]) < 0 && NextI < ArgNum &&
         getParentIndex(ArgTypes[NextI]) == I) {
       int64_t Alignment = getPartialStructRequiredAlignment(HstPtrBase);
-      Padding = (int64_t)HstPtrBegin % Alignment;
-      if (Padding) {
+      TgtPadding = (int64_t)HstPtrBegin % Alignment;
+      if (TgtPadding) {
         DP("Using a padding of %" PRId64 " bytes for begin address " DPxMOD
            "\n",
-           Padding, DPxPTR(HstPtrBegin));
-        HstPtrBegin = (char *)HstPtrBegin - Padding;
-        DataSize += Padding;
+           TgtPadding, DPxPTR(HstPtrBegin));
       }
     }
 
     // Address of pointer on the host and device, respectively.
     void *PointerHstPtrBegin, *PointerTgtPtrBegin;
     TargetPointerResultTy PointerTpr;
     bool IsHostPtr = false;
     bool IsImplicit = ArgTypes[I] & OMP_TGT_MAPTYPE_IMPLICIT;
if (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ) {
       // The map entry for s comes first, and the PTR_AND_OBJ entry comes
       // afterward, so the pointer is already allocated by the time the
       // PTR_AND_OBJ entry is handled below, and PointerTgtPtrBegin is thus
       // non-null.  However, "declare target link" can produce a PTR_AND_OBJ
       // entry for a global that might not already be allocated by the time the
       // PTR_AND_OBJ entry is handled below, and so the allocation might fail
       // when HasPresentModifier.
       PointerTpr = Device.getTargetPointer(
-          HDTTMap, HstPtrBase, HstPtrBase, sizeof(void *),
+          HDTTMap, HstPtrBase, HstPtrBase, /*TgtPadding=*/0, sizeof(void *),
           /*HstPtrName=*/nullptr,
           /*HasFlagTo=*/false, /*HasFlagAlways=*/false, IsImplicit, UpdateRef,
           HasCloseModifier, HasPresentModifier, HasHoldModifier, AsyncInfo,
           /* OwnedTPR */ nullptr, /* ReleaseHDTTMap */ false);
       PointerTgtPtrBegin = PointerTpr.TargetPointer;
       IsHostPtr = PointerTpr.Flags.IsHostPointer;
       if (!PointerTgtPtrBegin) {
         REPORT("Call to getTargetPointer returned null pointer (%s).\n",
if (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ) {
       UpdateRef =
           (!FromMapper || I != 0); // subsequently update ref count of pointee
     }
 
     const bool HasFlagTo = ArgTypes[I] & OMP_TGT_MAPTYPE_TO;
     const bool HasFlagAlways = ArgTypes[I] & OMP_TGT_MAPTYPE_ALWAYS;
     // Note that HDTTMap will be released in getTargetPointer.
     auto TPR = Device.getTargetPointer(
-        HDTTMap, HstPtrBegin, HstPtrBase, DataSize, HstPtrName, HasFlagTo,
-        HasFlagAlways, IsImplicit, UpdateRef, HasCloseModifier,
+        HDTTMap, HstPtrBegin, HstPtrBase, TgtPadding, DataSize, HstPtrName,
+        HasFlagTo, HasFlagAlways, IsImplicit, UpdateRef, HasCloseModifier,
         HasPresentModifier, HasHoldModifier, AsyncInfo, PointerTpr.getEntry());
     void *TgtPtrBegin = TPR.TargetPointer;
     IsHostPtr = TPR.Flags.IsHostPointer;
     // If data_size==0, then the argument could be a zero-length pointer to
     // NULL, so getOrAlloc() returning NULL is not an error.
     if (!TgtPtrBegin && (DataSize || HasPresentModifier)) {
       REPORT("Call to getTargetPointer returned null pointer (%s).\n",
              HasPresentModifier ? "'present' map type modifier"
if (ArgMappers && ArgMappers[I]) {
         return OFFLOAD_FAIL;
       }
 
       // Skip the rest of this function, continue to the next argument.
       continue;
     }
 
     void *HstPtrBegin = Args[I];
-    void *HstPtrBase = ArgBases[I];
     int64_t DataSize = ArgSizes[I];
-    // Adjust for proper alignment if this is a combined entry (for structs).
-    // Look at the next argument - if that is MEMBER_OF this one, then this one
-    // is a combined entry.
-    const int NextI = I + 1;
-    if (getParentIndex(ArgTypes[I]) < 0 && NextI < ArgNum &&
-        getParentIndex(ArgTypes[NextI]) == I) {
-      int64_t Alignment = getPartialStructRequiredAlignment(HstPtrBase);
-      int64_t Padding = (int64_t)HstPtrBegin % Alignment;
-      if (Padding) {
-        DP("Using a Padding of %" PRId64 " bytes for begin address " DPxMOD
-           "\n",
-           Padding, DPxPTR(HstPtrBegin));
-        HstPtrBegin = (char *)HstPtrBegin - Padding;
-        DataSize += Padding;
-      }
-    }
-
     bool IsImplicit = ArgTypes[I] & OMP_TGT_MAPTYPE_IMPLICIT;
     bool UpdateRef = (!(ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF) ||
                       (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ)) &&
                      !(FromMapper && I == 0);
     bool ForceDelete = ArgTypes[I] & OMP_TGT_MAPTYPE_DELETE;
     bool HasPresentModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_PRESENT;
     bool HasHoldModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_OMPX_HOLD;
 

openmp/libomptarget/test/mapping/padding_not_mapped.c

+// RUN: %libomptarget-compile-generic -fopenmp-version=51
+// RUN: %libomptarget-run-fail-generic 2>&1 \
+// RUN: | %fcheck-generic
+
+// The host memory layout for the following program looks like this:
+//
+//   | 4 bytes | 4 bytes | 8 bytes |
+//   |   s.x   |   s.y   |   s.z   |
+//   `-----------------------------'
+//
+// s is always at least 8-byte aligned in host memory due to s.z, so
+// libomptarget's device padding for map(s.y,s.z) always maps to host memory
+// that includes s.x.  At one time, s.x appeared to be mapped as a result, but
+// libomptarget has since been fixed not to consider device padding as mapped to
+// host memory.
+
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+  struct S { int x; int y; double z; } s = {1, 2, 3};
+
+  // CHECK: &s.x = 0x[[#%x,HOST_ADDR:]], size = [[#%u,SIZE:]]
+  fprintf(stderr, "&s = %p\n", &s);
+  fprintf(stderr, "&s.x = %p, size = %ld\n", &s.x, sizeof s.x);
+  fprintf(stderr, "&s.y = %p\n", &s.y);
+  fprintf(stderr, "&s.z = %p\n", &s.z);
+
+  // CHECK: s.x is present: 0
+  // CHECK: s.x = 1{{$}}
+  #pragma omp target enter data map(alloc: s.y, s.z)
+  int dev = omp_get_default_device();
+  fprintf(stderr, "s.x is present: %d\n", omp_target_is_present(&s.x, dev));
+  #pragma omp target update from(s.x) // should have no effect
+  fprintf(stderr, "s.x = %d\n", s.x);
+
+  // CHECK: Libomptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]] bytes)
+  // CHECK: Libomptarget error: Call to getTargetPointer returned null pointer ('present' map type modifier).
+  // CHECK: Libomptarget fatal error 1: failure of target construct while offloading is mandatory
+  #pragma omp target enter data map(present, alloc: s.x)
+
+  return 0;
+}

openmp/libomptarget/test/mapping/power_of_two_alignment.c

 //   padding of that size for s, but at most 1 byte is ever actually needed.
 //
 // Below, we try many sizes of stackPad to try to produce those conditions.
 //
 // When collidePost was then mapped to the same host memory as the unnecessary
 // padding for s, libomptarget reported an array extension error.  collidePost
 // is never fully contained within that padding (which would avoid the extension
 // error) because collidePost is 16 bytes while the padding is always less than
-// 16 bytes due to the modulo operations.
+// 16 bytes due to the modulo operations.  (Later, libomptarget was changed not
+// to consider padding to be mapped to the host, so it cannot be involved in
+// array extension errors.)
 
 #include <stdint.h>
 #include <stdio.h>
 
 template <typename StackPad>
 void test() {
   StackPad stackPad;
   struct S { char x; char y[7]; char z[8]; };