[Libomptarget] Implement OpenMP 5.2 semantics for device pointers

[jhuber6]

Summary

In OpenMP 5.2, §5.8.6, page 160 line 32-33, when a device pointer
allocated by omp_target_alloc has implicitly been included on a target
construct as a zero-length array, the pointer initialisation should not
find a matching mapped list item, and so should retain its value as a
firstprivate variable. Previously, we would return a null pointer if the
list item was not found. This patch updates the map handling to the
OpenMP 5.2 semantics.

[jdoerfert]

LG

[ye-luo]

openmp/libomptarget/include/device.h

285

IsPresent is not orthogonal to IsHostPointer. This unamed struct design is flawed.

Also directly access TargetPointerResultTy member is also bad. There is lack of read only protection. accessing members should via a function.

something like.
struct TargetPointerResultTy {

std::unique_ptr<EntryFlags> flags# null is the entry doesn't exist.

inline bool isPresent() const { return bool(flags); }

}

openmp/libomptarget/test/mapping/implicit_device_ptr.c

13

Doesn't compile. Need (int *) omp_target_alloc

[jhuber6]

openmp/libomptarget/include/device.h

285

IsPresent is not orthogonal to IsHostPointer. This unamed struct design is flawed.

Could you elaborate on this? They should mean different things, with the current semantics we pass the pointer first-private, but do not assume it is a valid host pointer. That's why we use both IsPresent and IsHostPointer. For USM we will have IsPresent be false and IsHostPointer be true, for this implicit mapping IsHostPointer will be false.

Also directly access TargetPointerResultTy member is also bad. There is lack of read only protection. accessing members should via a function.

Sure, I'll add it. I'm not exactly sure what the utility of this struct is anyway, considering that these are just bools.

openmp/libomptarget/test/mapping/implicit_device_ptr.c

13

This is C code, the cast shouldn't be necessary.

[jhuber6]

Updating according to @ye-luo's suggestions.

[jdoerfert]

openmp/libomptarget/include/device.h

285

1. IsPresent and IsHostPointer are orthogonal, as elaborated this morning in the call.
2. All the rest is unrelated to this patch. Potentially fair but unrelated.

[ye-luo]

openmp/libomptarget/include/device.h

285

OK. I was thinking of IsHostPointer for zero-length array. Actually your are right they are for USM. There is still some interaction between isNew and isPresent. I think isNew only make sense when isPresent is true.

Maybe std::unque_ptr is not a good option because it involves memory alloc/dealloc, std::optional is better but it requires c++17. Probably still need to stay in C++14 now. So the current way is acceptable.

openmp/libomptarget/src/device.cpp

433

I'm wondering if the USM case and this case should be fused in a more compact way.
The action is the same, when a pointer is not present in the map table. Keep the original value.

Even the logic of USM is fishy https://github.com/llvm/llvm-project/blob/178554f3c8f983bd192818b6d46e4d95560c4964/openmp/libomptarget/src/api.cpp#L133
But I would consider doing any related change in a separate patch.

openmp/libomptarget/test/mapping/implicit_device_ptr.c

13

You are right, I was compiling as C++.

[ye-luo]

This comment was removed by ye-luo.

[ye-luo]

openmp/libomptarget/src/api.cpp

128

129-134 seems to be useless with this patch.

openmp/libomptarget/include/device.h

 /// 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;
     /// If the pointer is actually a host pointer (when unified memory enabled)
     unsigned IsHostPointer : 1;
-  } Flags = {0, 0};
+    /// If the pointer is present in the mapping table.
+    unsigned IsPresent : 1;
+  } Flags = {0, 0, 0};
+
+  bool isPresent() const { return Flags.IsPresent; }
+
+  bool isHostPointer() const { return Flags.IsHostPointer; }
 
   /// The corresponding map table entry which is stable.
   HostDataToTargetTy *Entry = nullptr;
 
   /// The corresponding target pointer
   void *TargetPointer = nullptr;
 };
 

openmp/libomptarget/src/api.cpp

EXTERN int omp_target_is_present(const void *Ptr, int DeviceNum) {
   // omp_target_is_present tests whether a host pointer refers to storage that
   // is mapped to a given device. However, due to the lack of the storage size,
   // only check 1 byte. Cannot set size 0 which checks whether the pointer (zero
   // lengh array) is mapped instead of the referred storage.
   TargetPointerResultTy TPR =
       Device.getTgtPtrBegin(const_cast<void *>(Ptr), 1, 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;
+  int Rc = TPR.isPresent();
   DP("Call to omp_target_is_present returns %d\n", Rc);
   return Rc;
 }
 
 EXTERN int omp_target_memcpy(void *Dst, const void *Src, size_t Length,
                              size_t DstOffset, size_t SrcOffset, int DstDevice,
                              int SrcDevice) {
   TIMESCOPE();

openmp/libomptarget/src/device.cpp

TargetPointerResultTy DeviceTy::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) {
   HDTTMapAccessorTy HDTTMap = HostDataToTargetMap.getExclusiveAccessor();
 
   void *TargetPointer = nullptr;
   bool IsHostPtr = false;
+  bool IsPresent = true;
   bool IsNew = false;
 
   LookupResult LR = lookupMapping(HDTTMap, 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
if (LR.Flags.IsContained ||
     // 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.
     if (Size) {
       DP("Return HstPtrBegin " DPxMOD " Size=%" PRId64 " for unified shared "
          "memory\n",
          DPxPTR((uintptr_t)HstPtrBegin), Size);
+      IsPresent = false;
       IsHostPtr = true;
       TargetPointer = HstPtrBegin;
     }
   } else if (HasPresentModifier) {
     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 "
if (LR.Flags.IsContained ||
     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,
          Entry->dynRefCountToStr().c_str(), Entry->holdRefCountToStr().c_str(),
          (HstPtrName) ? getNameFromMapping(HstPtrName).c_str() : "unknown");
     TargetPointer = (void *)Ptr;
+  } else {
+    // This entry is not present and we did not create a new entry for it.
+    IsPresent = false;
   }
 
   // If the target pointer is valid, and we need to transfer data, issue the
   // data transfer.
   if (TargetPointer && !IsHostPtr && HasFlagTo && (IsNew || HasFlagAlways)) {
     // Lock the entry before releasing the mapping table lock such that another
     // thread that could issue data movement will get the right result.
     std::lock_guard<decltype(*Entry)> LG(*Entry);
if (Entry && !IsHostPtr && !HasPresentModifier) {
           return {{false /* IsNewEntry */, false /* IsHostPointer */},
                   nullptr /* Entry */,
                   nullptr /* TargetPointer */};
         }
       }
     }
   }
 
-  return {{IsNew, IsHostPtr}, Entry, TargetPointer};
+  return {{IsNew, IsHostPtr, IsPresent}, 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) {
   HDTTMapAccessorTy HDTTMap = HostDataToTargetMap.getExclusiveAccessor();
 
   void *TargetPointer = NULL;
   bool IsNew = false;
+  bool IsPresent = true;
   IsHostPtr = false;
   IsLast = false;
   LookupResult LR = lookupMapping(HDTTMap, HstPtrBegin, Size);
 
   if (LR.Flags.IsContained ||
       (!MustContain && (LR.Flags.ExtendsBefore || LR.Flags.ExtendsAfter))) {
     auto &HT = *LR.Entry;
     IsLast = HT.decShouldRemove(UseHoldRefCount, ForceDelete);
if (LR.Flags.IsContained ||
     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
     // is on then it means we have stumbled upon a value which we need to
     // use directly from the host.
     DP("Get HstPtrBegin " DPxMOD " Size=%" PRId64 " for unified shared "
        "memory\n",
        DPxPTR((uintptr_t)HstPtrBegin), Size);
+    IsPresent = false;
     IsHostPtr = true;
     TargetPointer = HstPtrBegin;
+  } else {
+    // OpenMP Specification v5.2: if a matching list item is not found, the
+    // pointer retains its original value as per firstprivate semantics.
+    IsPresent = false;
+    IsHostPtr = false;
+    TargetPointer = HstPtrBegin;
   }
 
-  return {{IsNew, IsHostPtr}, LR.Entry, TargetPointer};
+  return {{IsNew, IsHostPtr, IsPresent}, LR.Entry, TargetPointer};
 }
 
 // Return the target pointer begin (where the data will be moved).
 void *DeviceTy::getTgtPtrBegin(HDTTMapAccessorTy &HDTTMap, void *HstPtrBegin,
                                int64_t Size) {
   uintptr_t HP = (uintptr_t)HstPtrBegin;
   LookupResult LR = lookupMapping(HDTTMap, HstPtrBegin, Size);
   if (LR.Flags.IsContained || LR.Flags.ExtendsBefore || LR.Flags.ExtendsAfter) {

openmp/libomptarget/src/omptarget.cpp

for (int32_t I = ArgNum - 1; I >= 0; --I) {
     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);
     void *TgtPtrBegin = TPR.TargetPointer;
-    if (!TgtPtrBegin && (DataSize || HasPresentModifier)) {
+    if (!TPR.isPresent() && !TPR.isHostPointer() &&
+        (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
         // on entry to the region if the corresponding list item does not appear
         // in the device data environment then an error occurs and the program
if (!TPR.isPresent() && !TPR.isHostPointer() &&
          DataSize, DPxPTR(TgtPtrBegin), (IsLast ? "" : " not"));
     }
 
     // OpenMP 5.1, sec. 2.21.7.1 "map Clause", p. 351 L14-16:
     // "If the map clause appears on a target, target data, or target exit data
     // construct and a corresponding list item of the original list item is not
     // present in the device data environment on exit from the region then the
     // list item is ignored."
-    if (!TgtPtrBegin)
+    if (!TPR.isPresent())
       continue;
 
     bool DelEntry = IsLast;
 
     // If the last element from the mapper (for end transfer args comes in
     // reverse order), do not remove the partial entry, the parent struct still
     // exists.
     if ((ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF) &&
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,
       /*UseHoldRefCount=*/false, IsHostPtr, /*MustContain=*/true);
   void *TgtPtrBegin = TPR.TargetPointer;
-  if (!TgtPtrBegin) {
+  if (!TPR.isPresent()) {
     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);
       return OFFLOAD_FAIL;
     }
     return OFFLOAD_SUCCESS;
if (!(ArgTypes[I] & OMP_TGT_MAPTYPE_TARGET_PARAM)) {
         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);
         PointerTgtPtrBegin = TPR.TargetPointer;
-        if (!PointerTgtPtrBegin) {
+        if (!TPR.isPresent()) {
           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"
              "variable (" DPxMOD ")\n",
              DPxPTR(HstPtrVal));

openmp/libomptarget/test/mapping/implicit_device_ptr.c

+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <omp.h>
+#include <stdio.h>
+
+// OpenMP 5.1. sec 5.8.6 "Pointer Initialization for Device Data Environments"
+// p. 160 L32-33: "If a matching mapped list item is not found, the pointer
+// retains its original value as per the32 firstprivate semantics described in
+// Section 5.4.4."
+
+int main(void) {
+  int *A = (int *)omp_target_alloc(sizeof(int), omp_get_default_device());
+
+#pragma omp target
+  { *A = 1; }
+
+  int Result = 0;
+#pragma omp target map(from : Result)
+  { Result = *A; }
+
+  // CHECK: PASS
+  if (Result == 1)
+    printf("PASS\n");
+
+  omp_target_free(A, omp_get_default_device());
+}