[libomptarget] Correctly return the implicit device base address when the base itself has not been mapped.

[grokos]

Summary

When an array is mapped starting from an offset, its base address does not fall within the mapped range. If the same array is later captured implicitly in a target region or we use a use_device_ptr clause, the array appears as unmapped. This leads to otherwise legal OpenMP programs failing to execute correctly.

[ye-luo]

openmp/libomptarget/test/mapping/array_section_implicit_capture.c

32

When two distinct sections of A is mapped, implicit map looks up the base address, then which section to pick.

I think in this example, the implicit map should fail to look up A.
Users needs to explicitly add map(to: A[FROM:LENGTH], B[FROM:LENGTH]) to be valid.

[grokos]

openmp/libomptarget/test/mapping/array_section_implicit_capture.c

32

We only map one section, mapping two distinct sections of the same object is illegal. Implicitly mapping more, i.e. asking for the base when only a section in the middle of the object has been mapped is allowed. This example is valid OpenMP usage.

[ABataev]

openmp/libomptarget/test/mapping/array_section_implicit_capture.c

32

IIRC, Ye is correct here. Mapping 2 sections in different constructs is legal, just the inner one must be a subset of the outer.

[grokos]

openmp/libomptarget/test/mapping/array_section_implicit_capture.c

32

Sure, re-mapping what is already mapped is perfectly fine. I'm talking about implicitly "mapping" more, i.e. not really mapping a superset but rather using a zero-length pointer (or the whole named array) to refer to the object just for the sake of convenience (without having to re-write the exact range). Libomptarget has always supported this scenario, e.g. we have mapped A[0 : N/2] and then pass A[0:0] (or A[0:N] if A is a named array) to the target region. Now, if we split the array across multiple devices and each device only maps a chunk, without this patch all devices (except for the first one) will fail to look up the mapping.

[ye-luo]

openmp/libomptarget/test/mapping/array_section_implicit_capture.c

32

I think the spec has taken into account what I have mentioned and there is 5.8.6 Pointer Initialization for Device Data Environments in the spec. Clearly, base pointer is the way to go.

[grokos]

Added test for use_device_ptr

[grokos]

Changed the title because it was totally misleading.

[RaviNarayanaswamy]

openmp/libomptarget/test/mapping/array_section_implicit_capture.c

32

When doing look up cannot have same base for 2 different sections of the array
See
n
17 2. The corresponding pointer variable becomes an attached pointer for the corresponding list item.
18 3. If the original base pointer and the corresponding attached pointer share storage, then the
19 original list item and the corresponding list item must share storag

[RaviNarayanaswamy]

openmp/libomptarget/test/mapping/array_section_implicit_capture.c

32

Ye, are you asking if there are 2 sections mapped form same array ie they have same base then the look up should fail?

[ye-luo]

openmp/libomptarget/test/mapping/array_section_implicit_capture.c

32

Ravi, I was thinking of

#include <iostream>

int main()
{
  int a[30];
  a[ 1] = 0;
  a[21] = 0;
  #pragma omp target map(from: a[:10], a[20:10])
  {
    a[ 1] = 1;
    a[21] = 2;
  }
  std::cout << a[1] << " " << a[21] << std::endl;
}

But probably this code is illegal.

[RaviNarayanaswamy]

openmp/libomptarget/test/mapping/array_section_implicit_capture.c

32

#pragma omp target map(from: a[:10], a[20:10])  is illegal

[grokos]

Added a check to catch the error when the user tries to map distinct sections of the same object using different pointers.

openmp/libomptarget/include/device.h

 
 typedef std::set<HostDataToTargetTy, std::less<>> HostDataToTargetListTy;
 
 struct LookupResult {
   struct {
     unsigned IsContained : 1;
     unsigned ExtendsBefore : 1;
     unsigned ExtendsAfter : 1;
+    unsigned OnlyBaseFound : 1;
   } Flags;
 
   HostDataToTargetListTy::iterator Entry;
 
-  LookupResult() : Flags({0, 0, 0}), Entry() {}
+  LookupResult() : Flags({0, 0, 0, 0}), Entry() {}
 };
 
 /// This struct will be returned by \p DeviceTy::getTargetPointer which provides
 /// more data than just a target pointer.
 struct TargetPointerResultTy {
   struct {
     /// If the map table entry is just created
     unsigned IsNewEntry : 1;
struct DeviceTy {
   DeviceTy(const DeviceTy &D) = delete;
   DeviceTy &operator=(const DeviceTy &D) = delete;
 
   ~DeviceTy();
 
   // Return true if data can be copied to DstDevice directly
   bool isDataExchangable(const DeviceTy &DstDevice);
 
-  LookupResult lookupMapping(void *HstPtrBegin, int64_t Size);
+  LookupResult lookupMapping(void *HstPtrBase, void *HstPtrBegin, int64_t Size);
   /// Get the target pointer based on host pointer begin and base. If the
   /// mapping already exists, the target pointer will be returned directly. In
   /// addition, if required, the memory region pointed by \p HstPtrBegin of size
   /// \p Size will also be transferred to the device. If the mapping doesn't
   /// 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(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);
   void *getTgtPtrBegin(void *HstPtrBegin, int64_t Size);
-  TargetPointerResultTy getTgtPtrBegin(void *HstPtrBegin, int64_t Size,
-                                       bool &IsLast, bool UpdateRefCount,
+  TargetPointerResultTy getTgtPtrBegin(void *HstPtrBegin, void *HstPtrBase,
+                                       int64_t Size, bool &IsLast,
+                                       bool UpdateRefCount,
                                        bool UseHoldRefCount, bool &IsHostPtr,
                                        bool MustContain = false,
                                        bool ForceDelete = false);
   /// For the map entry for \p HstPtrBegin, decrement the reference count
   /// specified by \p HasHoldModifier and, if the the total reference count is
   /// then zero, deallocate the corresponding device storage and remove the map
   /// entry.  Return \c OFFLOAD_SUCCESS if the map entry existed, and return
   /// \c OFFLOAD_FAIL if not.  It is the caller's responsibility to skip calling

openmp/libomptarget/src/api.cpp

DP("Call to omp_target_is_present with invalid device ID, returning "
        "false\n");
     return false;
   }
 
   DeviceTy &Device = *PM->Devices[device_num];
   bool IsLast; // not used
   bool IsHostPtr;
   TargetPointerResultTy TPR =
-      Device.getTgtPtrBegin(const_cast<void *>(ptr), 0, IsLast,
-                            /*UpdateRefCount=*/false,
+      Device.getTgtPtrBegin(const_cast<void *>(ptr), const_cast<void *>(ptr), 0,
+                            IsLast, /*UpdateRefCount=*/false,
                             /*UseHoldRefCount=*/false, IsHostPtr);
   int rc = (TPR.TargetPointer != NULL);
   // Under unified memory the host pointer can be returned by the
   // getTgtPtrBegin() function which means that there is no device
   // corresponding point for ptr. This function should return false
   // in that situation.
   if (PM->RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY)
     rc = !IsHostPtr;

openmp/libomptarget/src/device.cpp

if (search != HostDataToTargetMap.end()) {
     REPORT("Association not found\n");
   }
 
   // Mapping not found
   DataMapMtx.unlock();
   return OFFLOAD_FAIL;
 }
 
-LookupResult DeviceTy::lookupMapping(void *HstPtrBegin, int64_t Size) {
+LookupResult DeviceTy::lookupMapping(void *HstPtrBase, void *HstPtrBegin,
+                                     int64_t Size) {
   uintptr_t hp = (uintptr_t)HstPtrBegin;
+  uintptr_t hb = (uintptr_t)HstPtrBase;
   LookupResult lr;
 
   DP("Looking up mapping(HstPtrBegin=" DPxMOD ", Size=%" PRId64 ")...\n",
      DPxPTR(hp), Size);
 
   if (HostDataToTargetMap.empty())
     return lr;
 
   auto upper = HostDataToTargetMap.upper_bound(hp);
   // check the left bin
   if (upper != HostDataToTargetMap.begin()) {
     lr.Entry = std::prev(upper);
     auto &HT = *lr.Entry;
     // Is it contained?
     lr.Flags.IsContained = hp >= HT.HstPtrBegin && hp < HT.HstPtrEnd &&
                            (hp + Size) <= HT.HstPtrEnd;
     // Does it extend beyond the mapped region?
     lr.Flags.ExtendsAfter = hp < HT.HstPtrEnd && (hp + Size) > HT.HstPtrEnd;
+    // Did we just find the base?
+    // There is a special case where we map a struct via a double pointer, e.g.
+    // typedef struct {...} S;
+    // S s;
+    // S *sp = &s;
+    // S **spp = &sp;
+    // #pragma omp target map (to: spp[0][0]) {...}
+    // In that case, the record for pointer sp will share the same base as the
+    // record for struct s, but we shouldn't report that we only found the base,
+    // i.e. we must let libomptarget create a new record in HostDataToTargetMap
+    // for the struct s itself which will have the same base as the record for
+    // pointer sp. This scenario can only happen in the left bin.
+    if (!lr.Flags.IsContained && !lr.Flags.ExtendsAfter &&
+        hb == HT.HstPtrBase &&
+        HT.HstPtrEnd - HT.HstPtrBegin != sizeof(void *))
+      lr.Flags.OnlyBaseFound = true;
   }
 
   // check the right bin
-  if (!(lr.Flags.IsContained || lr.Flags.ExtendsAfter) &&
-      upper != HostDataToTargetMap.end()) {
+  if (!(lr.Flags.IsContained || lr.Flags.ExtendsAfter || lr.Flags.OnlyBaseFound)
+      && upper != HostDataToTargetMap.end()) {
     lr.Entry = upper;
     auto &HT = *lr.Entry;
     // Does it extend into an already mapped region?
     lr.Flags.ExtendsBefore =
         hp < HT.HstPtrBegin && (hp + Size) > HT.HstPtrBegin;
     // Does it extend beyond the mapped region?
     lr.Flags.ExtendsAfter = hp < HT.HstPtrEnd && (hp + Size) > HT.HstPtrEnd;
+    // Did we just find the base?
+    if (!lr.Flags.ExtendsBefore && !lr.Flags.ExtendsAfter &&
+        hb == HT.HstPtrBase)
+      lr.Flags.OnlyBaseFound = true;
   }
 
   if (lr.Flags.ExtendsBefore) {
     DP("WARNING: Pointer is not mapped but section extends into already "
        "mapped data\n");
   }
   if (lr.Flags.ExtendsAfter) {
     DP("WARNING: Pointer is already mapped but section extends beyond mapped "
DeviceTy::getTargetPointer(void *HstPtrBegin, void *HstPtrBase, int64_t Size,
                            bool HasPresentModifier, bool HasHoldModifier,
                            AsyncInfoTy &AsyncInfo) {
   void *TargetPointer = nullptr;
   bool IsHostPtr = false;
   bool IsNew = false;
 
   DataMapMtx.lock();
 
-  LookupResult LR = lookupMapping(HstPtrBegin, Size);
+  LookupResult LR = lookupMapping(HstPtrBase, HstPtrBegin, Size);
   auto Entry = LR.Entry;
 
   // Check if the pointer is contained.
   // If a variable is mapped to the device manually by the user - which would
   // lead to the IsContained flag to be true - then we must ensure that the
   // device address is returned even under unified memory conditions.
   if (LR.Flags.IsContained ||
-      ((LR.Flags.ExtendsBefore || LR.Flags.ExtendsAfter) && IsImplicit)) {
+      ((LR.Flags.ExtendsBefore || LR.Flags.ExtendsAfter) && IsImplicit) ||
+      (LR.Flags.OnlyBaseFound && !Size)) {
     auto &HT = *LR.Entry;
     const char *RefCountAction;
     assert(HT.getTotalRefCount() > 0 && "expected existing RefCount > 0");
     if (UpdateRefCount) {
       // After this, RefCount > 1.
       HT.incRefCount(HasHoldModifier);
       RefCountAction = " (incremented)";
     } else {
       // It might have been allocated with the parent, but it's still new.
       IsNew = HT.getTotalRefCount() == 1;
       RefCountAction = " (update suppressed)";
     }
     const char *DynRefCountAction = HasHoldModifier ? "" : RefCountAction;
     const char *HoldRefCountAction = HasHoldModifier ? RefCountAction : "";
-    uintptr_t Ptr = HT.TgtPtrBegin + ((uintptr_t)HstPtrBegin - HT.HstPtrBegin);
+
+    uintptr_t Ptr = 0;
+    if (LR.Flags.OnlyBaseFound)
+      // Return the implied device base
+      Ptr = HT.TgtPtrBegin - (HT.HstPtrBegin - HT.HstPtrBase);
+    else
+      // Return the device address corresponding to HstPtrBegin
+      Ptr = HT.TgtPtrBegin + ((uintptr_t)HstPtrBegin - HT.HstPtrBegin);
     INFO(OMP_INFOTYPE_MAPPING_EXISTS, DeviceID,
          "Mapping exists%s with HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD
          ", Size=%" PRId64 ", DynRefCount=%s%s, HoldRefCount=%s%s, Name=%s\n",
          (IsImplicit ? " (implicit)" : ""), DPxPTR(HstPtrBegin), DPxPTR(Ptr),
          Size, HT.dynRefCountToStr().c_str(), DynRefCountAction,
          HT.holdRefCountToStr().c_str(), HoldRefCountAction,
          (HstPtrName) ? getNameFromMapping(HstPtrName).c_str() : "unknown");
     TargetPointer = (void *)Ptr;
   } else if ((LR.Flags.ExtendsBefore || LR.Flags.ExtendsAfter) && !IsImplicit) {
     // Explicit extension of mapped data - not allowed.
     MESSAGE("explicit extension not allowed: host address specified is " DPxMOD
             " (%" PRId64
             " bytes), but device allocation maps to host at " DPxMOD
             " (%" PRId64 " bytes)",
             DPxPTR(HstPtrBegin), Size, DPxPTR(Entry->HstPtrBegin),
             Entry->HstPtrEnd - Entry->HstPtrBegin);
     if (HasPresentModifier)
       MESSAGE("device mapping required by 'present' map type modifier does not "
               "exist for host address " DPxMOD " (%" PRId64 " bytes)",
               DPxPTR(HstPtrBegin), Size);
+  } else if (LR.Flags.OnlyBaseFound && Size) {
+    // If we only found the base address but a size has been provided, it
+    // means that the user tried to map another distinct chunk of an already
+    // partially-mapped object, which is illegal.
+    MESSAGE("Found record of existing mapping with the requested base address "
+            DPxMOD " but disjoint mapped data, mapping of two distinct chunks "
+            "of the same object is not allowed.", DPxPTR(HstPtrBase));
   } else if (PM->RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY &&
              !HasCloseModifier) {
     // If unified shared memory is active, implicitly mapped variables that are
     // not privatized use host address. Any explicitly mapped variables also use
     // host address where correctness is not impeded. In all other cases maps
     // are respected.
     // In addition to the mapping rules above, the close map modifier forces the
     // mapping of the variable to the device.
DeviceTy::getTargetPointer(void *HstPtrBegin, void *HstPtrBase, int64_t Size,
 
   return {{IsNew, IsHostPtr}, Entry, TargetPointer};
 }
 
 // Used by targetDataBegin, targetDataEnd, targetDataUpdate and target.
 // Return the target pointer begin (where the data will be moved).
 // Decrement the reference counter if called from targetDataEnd.
 TargetPointerResultTy
-DeviceTy::getTgtPtrBegin(void *HstPtrBegin, int64_t Size, bool &IsLast,
-                         bool UpdateRefCount, bool UseHoldRefCount,
-                         bool &IsHostPtr, bool MustContain, bool ForceDelete) {
+DeviceTy::getTgtPtrBegin(void *HstPtrBegin, void *HstPtrBase, int64_t Size,
+                         bool &IsLast, bool UpdateRefCount,
+                         bool UseHoldRefCount, bool &IsHostPtr,
+                         bool MustContain, bool ForceDelete) {
   void *TargetPointer = NULL;
   bool IsNew = false;
   IsHostPtr = false;
   IsLast = false;
   DataMapMtx.lock();
-  LookupResult lr = lookupMapping(HstPtrBegin, Size);
+  LookupResult lr = lookupMapping(HstPtrBase, HstPtrBegin, Size);
 
   if (lr.Flags.IsContained ||
-      (!MustContain && (lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter))) {
+      (!MustContain && (lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter)) ||
+      lr.Flags.OnlyBaseFound) {
     auto &HT = *lr.Entry;
     // We do not zero the total reference count here.  deallocTgtPtr does that
     // atomically with removing the mapping.  Otherwise, before this thread
     // removed the mapping in deallocTgtPtr, another thread could retrieve the
     // mapping, increment and decrement back to zero, and then both threads
     // would try to remove the mapping, resulting in a double free.
     IsLast = HT.decShouldRemove(UseHoldRefCount, ForceDelete);
     const char *RefCountAction;
if (lr.Flags.IsContained ||
     } else if (IsLast) {
       RefCountAction = " (deferred final decrement)";
     } else {
       HT.decRefCount(UseHoldRefCount);
       RefCountAction = " (decremented)";
     }
     const char *DynRefCountAction = UseHoldRefCount ? "" : RefCountAction;
     const char *HoldRefCountAction = UseHoldRefCount ? RefCountAction : "";
-    uintptr_t tp = HT.TgtPtrBegin + ((uintptr_t)HstPtrBegin - HT.HstPtrBegin);
+
+    uintptr_t tp = 0;
+    if (lr.Flags.OnlyBaseFound)
+      // Return the implied device base
+      tp = HT.TgtPtrBegin - (HT.HstPtrBegin - HT.HstPtrBase);
+    else
+      // Return the device address corresponding to HstPtrBegin
+      tp = HT.TgtPtrBegin + ((uintptr_t)HstPtrBegin - HT.HstPtrBegin);
     INFO(OMP_INFOTYPE_MAPPING_EXISTS, DeviceID,
          "Mapping exists with HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD ", "
          "Size=%" PRId64 ", DynRefCount=%s%s, HoldRefCount=%s%s\n",
          DPxPTR(HstPtrBegin), DPxPTR(tp), Size, HT.dynRefCountToStr().c_str(),
          DynRefCountAction, HT.holdRefCountToStr().c_str(), HoldRefCountAction);
     TargetPointer = (void *)tp;
   } else if (PM->RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY) {
     // If the value isn't found in the mapping and unified shared memory
DeviceTy::getTgtPtrBegin(void *HstPtrBegin, void *HstPtrBase, int64_t Size,
   DataMapMtx.unlock();
   return {{IsNew, IsHostPtr}, lr.Entry, TargetPointer};
 }
 
 // Return the target pointer begin (where the data will be moved).
 // Lock-free version called when loading global symbols from the fat binary.
 void *DeviceTy::getTgtPtrBegin(void *HstPtrBegin, int64_t Size) {
   uintptr_t hp = (uintptr_t)HstPtrBegin;
-  LookupResult lr = lookupMapping(HstPtrBegin, Size);
+  LookupResult lr = lookupMapping(nullptr, HstPtrBegin, Size);
   if (lr.Flags.IsContained || lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter) {
     auto &HT = *lr.Entry;
     uintptr_t tp = HT.TgtPtrBegin + (hp - HT.HstPtrBegin);
     return (void *)tp;
   }
 
   return NULL;
 }
 
 int DeviceTy::deallocTgtPtr(void *HstPtrBegin, int64_t Size,
                             bool HasHoldModifier) {
   // Check if the pointer is contained in any sub-nodes.
   int Ret = OFFLOAD_SUCCESS;
   DataMapMtx.lock();
-  LookupResult lr = lookupMapping(HstPtrBegin, Size);
+  LookupResult lr = lookupMapping(nullptr, HstPtrBegin, Size);
   if (lr.Flags.IsContained || lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter) {
     auto &HT = *lr.Entry;
     if (HT.decRefCount(HasHoldModifier) == 0) {
       DP("Deleting tgt data " DPxMOD " of size %" PRId64 "\n",
          DPxPTR(HT.TgtPtrBegin), Size);
       deleteData((void *)HT.TgtPtrBegin);
       INFO(OMP_INFOTYPE_MAPPING_CHANGED, DeviceID,
            "Removing map entry with HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD

 int32_t DeviceTy::deleteData(void *TgtPtrBegin) {
   return RTL->data_delete(RTLDeviceID, TgtPtrBegin);
 }
 
 // Submit data to device
 int32_t DeviceTy::submitData(void *TgtPtrBegin, void *HstPtrBegin, int64_t Size,
                              AsyncInfoTy &AsyncInfo) {
   if (getInfoLevel() & OMP_INFOTYPE_DATA_TRANSFER) {
-    LookupResult LR = lookupMapping(HstPtrBegin, Size);
+    LookupResult LR = lookupMapping(nullptr, HstPtrBegin, Size);
     auto *HT = &*LR.Entry;
 
     INFO(OMP_INFOTYPE_DATA_TRANSFER, DeviceID,
          "Copying data from host to device, HstPtr=" DPxMOD ", TgtPtr=" DPxMOD
          ", Size=%" PRId64 ", Name=%s\n",
          DPxPTR(HstPtrBegin), DPxPTR(TgtPtrBegin), Size,
          (HT && HT->HstPtrName) ? getNameFromMapping(HT->HstPtrName).c_str()
                                 : "unknown");
   }
 
   if (!AsyncInfo || !RTL->data_submit_async || !RTL->synchronize)
     return RTL->data_submit(RTLDeviceID, TgtPtrBegin, HstPtrBegin, Size);
   else
     return RTL->data_submit_async(RTLDeviceID, TgtPtrBegin, HstPtrBegin, Size,
                                   AsyncInfo);
 }
 
 // Retrieve data from device
 int32_t DeviceTy::retrieveData(void *HstPtrBegin, void *TgtPtrBegin,
                                int64_t Size, AsyncInfoTy &AsyncInfo) {
   if (getInfoLevel() & OMP_INFOTYPE_DATA_TRANSFER) {
-    LookupResult LR = lookupMapping(HstPtrBegin, Size);
+    LookupResult LR = lookupMapping(nullptr, HstPtrBegin, Size);
     auto *HT = &*LR.Entry;
     INFO(OMP_INFOTYPE_DATA_TRANSFER, DeviceID,
          "Copying data from device to host, TgtPtr=" DPxMOD ", HstPtr=" DPxMOD
          ", Size=%" PRId64 ", Name=%s\n",
          DPxPTR(TgtPtrBegin), DPxPTR(HstPtrBegin), Size,
          (HT && HT->HstPtrName) ? getNameFromMapping(HT->HstPtrName).c_str()
                                 : "unknown");
   }

openmp/libomptarget/src/omptarget.cpp

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 Padding = (int64_t)HstPtrBegin % Alignment;
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;
 
     // If PTR_AND_OBJ, HstPtrBegin is address of pointee
     TargetPointerResultTy TPR = Device.getTgtPtrBegin(
-        HstPtrBegin, DataSize, IsLast, UpdateRef, HasHoldModifier, IsHostPtr,
-        !IsImplicit, ForceDelete);
+        HstPtrBegin, HstPtrBase, DataSize, IsLast, UpdateRef, HasHoldModifier,
+        IsHostPtr, !IsImplicit, ForceDelete);
     void *TgtPtrBegin = TPR.TargetPointer;
     if (!TgtPtrBegin && (DataSize || HasPresentModifier)) {
       DP("Mapping does not exist (%s)\n",
          (HasPresentModifier ? "'present' map type modifier" : "ignored"));
       if (HasPresentModifier) {
         // OpenMP 5.1, sec. 2.21.7.1 "map Clause", p. 350 L10-13:
         // "If a map clause appears on a target, target data, target enter data
         // or target exit data construct with a present map-type-modifier then

 }
 
 static int targetDataContiguous(ident_t *loc, DeviceTy &Device, void *ArgsBase,
                                 void *HstPtrBegin, int64_t ArgSize,
                                 int64_t ArgType, AsyncInfoTy &AsyncInfo) {
   TIMESCOPE_WITH_IDENT(loc);
   bool IsLast, IsHostPtr;
   TargetPointerResultTy TPR = Device.getTgtPtrBegin(
-      HstPtrBegin, ArgSize, IsLast, /*UpdateRefCount=*/false,
+      HstPtrBegin, ArgsBase, ArgSize, IsLast, /*UpdateRefCount=*/false,
       /*UseHoldRefCount=*/false, IsHostPtr, /*MustContain=*/true);
   void *TgtPtrBegin = TPR.TargetPointer;
   if (!TgtPtrBegin) {
     DP("hst data:" DPxMOD " not found, becomes a noop\n", DPxPTR(HstPtrBegin));
     if (ArgType & OMP_TGT_MAPTYPE_PRESENT) {
       MESSAGE("device mapping required by 'present' motion modifier does not "
               "exist for host address " DPxMOD " (%" PRId64 " bytes)",
               DPxPTR(HstPtrBegin), ArgSize);
if (!(ArgTypes[I] & OMP_TGT_MAPTYPE_TARGET_PARAM)) {
         bool IsLast, IsHostPtr; // IsLast is unused.
         void *TgtPtrBase =
             (void *)((intptr_t)TgtArgs[TgtIdx] + TgtOffsets[TgtIdx]);
         DP("Parent lambda base " DPxMOD "\n", DPxPTR(TgtPtrBase));
         uint64_t Delta = (uint64_t)HstPtrBegin - (uint64_t)HstPtrBase;
         void *TgtPtrBegin = (void *)((uintptr_t)TgtPtrBase + Delta);
         void *&PointerTgtPtrBegin = AsyncInfo.getVoidPtrLocation();
         TargetPointerResultTy TPR = Device.getTgtPtrBegin(
-            HstPtrVal, ArgSizes[I], IsLast, /*UpdateRefCount=*/false,
-            /*UseHoldRefCount=*/false, IsHostPtr);
+            HstPtrVal, HstPtrBegin, ArgSizes[I], IsLast,
+            /*UpdateRefCount=*/false, /*UseHoldRefCount=*/false, IsHostPtr);
         PointerTgtPtrBegin = TPR.TargetPointer;
         if (!PointerTgtPtrBegin) {
           DP("No lambda captured variable mapped (" DPxMOD ") - ignored\n",
              DPxPTR(HstPtrVal));
           continue;
         }
         if (IsHostPtr) {
           DP("Unified memory is active, no need to map lambda captured"
if (ArgTypes[I] & OMP_TGT_MAPTYPE_LITERAL) {
       if (Ret != OFFLOAD_SUCCESS) {
         REPORT("Failed to process %sprivate argument " DPxMOD "\n",
                (IsFirstPrivate ? "first-" : ""), DPxPTR(HstPtrBegin));
         return OFFLOAD_FAIL;
       }
     } else {
       if (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ)
         HstPtrBase = *reinterpret_cast<void **>(HstPtrBase);
-      TPR = Device.getTgtPtrBegin(HstPtrBegin, ArgSizes[I], IsLast,
+      TPR = Device.getTgtPtrBegin(HstPtrBegin, HstPtrBase, ArgSizes[I], IsLast,
                                   /*UpdateRefCount=*/false,
                                   /*UseHoldRefCount=*/false, IsHostPtr);
       TgtPtrBegin = TPR.TargetPointer;
       TgtBaseOffset = (intptr_t)HstPtrBase - (intptr_t)HstPtrBegin;
 #ifdef OMPTARGET_DEBUG
       void *TgtPtrBase = (void *)((intptr_t)TgtPtrBegin + TgtBaseOffset);
       DP("Obtained target argument " DPxMOD " from host pointer " DPxMOD "\n",
          DPxPTR(TgtPtrBase), DPxPTR(HstPtrBegin));

openmp/libomptarget/test/mapping/array_section_implicit_capture.c

+// RUN: %libomptarget-compile-generic -fopenmp-version=51
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic
+
+#include <stdio.h>
+#include <stdlib.h>
+
+#define N 1024
+#define FROM 64
+#define LENGTH 128
+
+int main()
+{
+  float *A = (float *) malloc(N * sizeof(float));
+  float *B = (float *) malloc(N * sizeof(float));
+  float *C = (float *) malloc(N * sizeof(float));
+
+  for (int i = 0; i < N; i++) {
+    C[i] = 0.0;
+  }
+
+  for (int i = 0; i < N; i++) {
+    A[i] = i;
+    B[i] = 2*i;
+  }
+
+  #pragma omp target enter data map(to: A[FROM:LENGTH], B[FROM:LENGTH])
+  #pragma omp target enter data map(alloc: C[FROM:LENGTH])
+
+  // A, B and C have been mapped starting at index FROM, but inside the kernel
+  // they are captured implicitly so the library must look them up using their
+  // base address.
+  #pragma omp target
+  {
+    for (int i = FROM; i < FROM+LENGTH; i++) {
+      C[i] = A[i] + B[i];
+    }
+  }
+
+  #pragma omp target exit data map(from: C[FROM:LENGTH])
+  #pragma omp target exit data map(delete: A[FROM:LENGTH], B[FROM:LENGTH])
+
+  int errors = 0;
+  for (int i = FROM; i < FROM+LENGTH; i++)
+    if (C[i] != A[i] + B[i])
+      ++errors;
+
+  // CHECK: Success
+  if (errors)
+    fprintf(stderr, "Failure\n");
+  else
+    fprintf(stderr, "Success\n");
+
+  free(A);
+  free(B);
+  free(C);
+
+  return 0;
+}

openmp/libomptarget/test/mapping/array_section_use_device_ptr.c

+// RUN: %libomptarget-compile-generic -fopenmp-version=51
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic
+
+#include <stdio.h>
+#include <stdlib.h>
+
+#define N 1024
+#define FROM 64
+#define LENGTH 128
+
+int main()
+{
+  float *A = (float *) malloc(N * sizeof(float));
+
+  #pragma omp target enter data map(to: A[FROM:LENGTH])
+
+  // A, has been mapped starting at index FROM, but inside the use_device_ptr
+  // clause it is captured by base so the library must look it up using the
+  // base address.
+
+  float *A_dev = NULL;
+  #pragma omp target data use_device_ptr(A)
+  {
+    A_dev = A;
+  }
+  #pragma omp target exit data map(delete: A[FROM:LENGTH])
+
+  // CHECK: Success
+  if (A_dev == NULL || A_dev == A)
+    fprintf(stderr, "Failure\n");
+  else
+    fprintf(stderr, "Success\n");
+
+  free(A);
+
+  return 0;
+}

openmp/libomptarget/test/mapping/two_chunks_via_ptrs.c

+// RUN: %libomptarget-compile-generic
+// RUN: %libomptarget-run-fail-generic 2>&1 \
+// RUN: | %fcheck-generic
+
+// CHECK: Libomptarget message: Found record of existing mapping with the requested base address 0x{{.*}} but disjoint mapped data, mapping of two distinct chunks of the same object is not allowed.
+// CHECK: Libomptarget error: Call to getTargetPointer returned null pointer (device failure or illegal mapping).
+// CHECK: Libomptarget fatal error 1: failure of target construct while offloading is mandatory
+
+#include <stdio.h>
+#include <stdlib.h>
+
+int main()
+{
+  int *mem = (int *) malloc(32 * sizeof(int));
+  int *a = mem;
+  int *b = mem;
+  #pragma omp target data map(from: a[0:10], b[20:10])
+  { }
+
+  return 0;
+}