[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

31

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

31

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

31

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

31

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

31

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

31

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

31

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

31

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

31

#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 {
   /// - 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 (!(lr.Flags.IsContained || lr.Flags.ExtendsAfter) &&
       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 looking up by base address, the entry can
+    // only be found in the right bin.
+    if (!lr.Flags.ExtendsBefore && !lr.Flags.ExtendsAfter &&
+        hp == 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,
   LookupResult LR = lookupMapping(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) {
     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;
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);
 
   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

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], C[FROM:LENGTH])
+
+  // A and B 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])
+
+  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;
+}