[OpenMP] Implement TR8 `present` map type modifier in runtime (2/2)

[jdenny]

Summary

This implements OpenMP runtime support for the OpenMP TR8 present
map type modifier. The previous patch in this series implements Clang
front end support. See that patch summary for behaviors that are not
yet supported.

[grokos]

openmp/libomptarget/src/device.cpp

209–212

I think this else-if should be moved right after the else-if that checks for explicit extension of mapping (i.e. after line 194) and outside the else if (Size) branch. If we have the present modifier then the data must be mapped already no matter whether a size is specified.

openmp/libomptarget/src/omptarget.cpp

320–336

Correct. The only case we can have a PTR_AND_OBJ entry is if we have a pointer which is member of a struct or element of an array. This means that whenever we have something like:

#pragma omp target map(s.p[0:N])

clang will first map the struct/array and then we'll visit the PTR_AND_OBJ entry, so the space for the pointer itself has been allocated already. I think it's safe to skip this check here (or you can leave it as a form of defensive programming).

[jdoerfert]

openmp/libomptarget/src/omptarget.cpp

492

Can we print the error in non debug mode as well? Is there a way we can provide more information on the location and variable?

[jdenny]

openmp/libomptarget/src/device.cpp

209–212

In the case of Size == 0, as far as I can tell, nothing is ever placed in HostDataToTargetMap, with or without the patch under review. Instead, getOrAllocTgtPtr always handles this case as already mapped and returns NULL for it. I believe your suggested change would not affect this behavior. However, when HasPresentModifier, your suggested change would cause the Mapping required but does not exist debug message to be printed for this case, incorrectly implying that this case is never already mapped and that a runtime error should follow. A runtime error won't follow because Size == 0.

In the case of Size != 0, your suggested change would affect unified shared memory. As far as I can tell, unified shared memory is handled similarly to the case of Size == 0: it's never placed in HostDataToTargetMap, and getOrAllocTgtPtr always handles it as already mapped. However, when HasPresentModifier, your suggested change would cause getOrAllocTgtPtr to always handle unified shared memory as unmapped, and a runtime error will follow, so unified shared memory is then unusable with the present modifier.

In both cases, your suggested change seems to produce the wrong behavior. Do you agree?

I need to add tests to cover these cases.

openmp/libomptarget/src/omptarget.cpp

320–336

Thanks for confirming. I'll leave it but adjust the comment not to be a todo.

483–498

This should be TgtPtrBegin || !data_size, mirroring similar code in target_data_begin.

492

I would prefer that too. When I investigated this originally, I didn't find a good precedent in this runtime implementation, so I decided to be conservative. Is there an existing diagnostic I should model this after, or is this a unique case?

[jdenny]

Added tests for cases of size-zero mappings and unified shared memory.

Fixed handling of size-zero mapping in target_data_end.

Converted TODO into a normal comment.

[jdenny]

openmp/libomptarget/src/device.h

183

I see no remaining getOrAllocTgtPtr calls that make use of these default arguments. Any objection to removing them in this patch?

[grokos]

openmp/libomptarget/src/device.cpp

209–212

In the case of Size == 0, as far as I can tell, nothing is ever placed in HostDataToTargetMap, with or without the patch under review. Instead, getOrAllocTgtPtr always handles this case as already mapped and returns NULL for it.

This is not always true. We can have zero size and still get a valid device address. Consider the example below:

int *p = malloc(N);
#pragma omp target enter data map(to: p[0:N])
...
#pragma omp target
{
  p[0] = 1;
}

Here p[0:N] is mapped via enter data and address &p[0] is inserted into HostDataToTargetMap. When we encounter the target region, the compiler captures pointer p implicitly as it appears inside the target region, however the compiler doesn't know its size (it's a pointer, not a named array), therefore it will be mapped as p[0:0], i.e. as a zero-size pointer. getOrAllocTgtPtr will find address &p[0] in HostDataToTargetMap (lr.isContained) and return the corresponding device address. If we have demanded that p[0:0] is present on the device, then libomptarget must check for its presence even if size is 0. So if you invoke the target region like this:

#pragma omp target map(present, alloc : p[0:0])

then libomptarget must either:

1. Return the corresponding device address of &p[0] if p[0:N] has been mapped before or
2. Error out if &p[0] cannot be found in HostDataToTargetMap (e.g. if we had omitted the enter data directive in the example above) but we requested it to be present

The latter case is not covered by this version of the patch.

Regarding unified shared memory you're right. So my proposed change should be:

if (contained) {
...
} else if (explicit extension) {
...
} else if (HasPresentModifier && (!USM || (USM && Close) ) {
  error out
} else if (Size) {
...
}

openmp/libomptarget/src/omptarget.cpp

492

Unfortunately, libomptarget does not have access to source code location. __tgt_ API functions do not have a related argument, like the ident_t *loc argument of __kmpc_ API functions in libomp. Of course, if there is demand for such a feature we can implement it, although we would need to change *EVERY* API function in libomptarget.

[jdenny]

openmp/libomptarget/src/device.cpp

209–212

Thanks for explaining. What is the correct behavior in the following case?

#pragma omp target data map(alloc : p[0:0])
#pragma omp target data map(present, alloc : p[0:0])

[jdenny]

Changed present modifier to fail for zero-length array section if the base address is not already mapped, and updated tests.

[jdenny]

openmp/libomptarget/src/device.cpp

209–212

Thanks for explaining. What is the correct behavior in the following case?

#pragma omp target data map(alloc : p[0:0])
#pragma omp target data map(present, alloc : p[0:0])

My latest update assumes the second directive should produce a runtime error because the first directive doesn't actually map anything. Please let me know if that's not right.

So my proposed change should be:

if (contained) {
...
} else if (explicit extension) {
...
} else if (HasPresentModifier && (!USM || (USM && Close) ) {

error out

} else if (Size) {
...
}

I organized this a little differently to avoid repeating the USM conditions, but I intend it to be equivalent.

[grokos]

openmp/libomptarget/src/device.cpp

209–212

Thanks for explaining. What is the correct behavior in the following case?

#pragma omp target data map(alloc : p[0:0])
#pragma omp target data map(present, alloc : p[0:0])

For the first one, if the object p is pointing to has been mapped before then there is no noticeable effect, the only thing that happens is that getOrAllocTgtPtr increments the RefCount of the mapping and returns the corresponding device address. If the object hasn't been mapped before and because size is zero, getOrAllocTgtPtr will return NULL (there is nothing to be allocated).

With the present modifier, if the object is already mapped, getOrAllocTgtPtr increments the RefCount and returns the corresponding device address (just like in the former case). If the object is not mapped, the present modifier will trigger an error - in other words NULL is not an option if we have requested something to be present.

openmp/libomptarget/src/device.h

183

Sure, they can go.

[jdenny]

openmp/libomptarget/src/device.cpp

209–212

Sorry, I meant that these directives are nested and the array was not previously mapped. It sounds like the outer directive does not map it in that case, and so the inner directive produces a runtime error.

[grokos]

openmp/libomptarget/src/device.cpp

209–212

Correct.

[jdenny]

Removed unused default arguments from getOrAllocTgtPtr.

[jdenny]

openmp/libomptarget/src/device.cpp

209–212

Thanks for all the careful explanations! I believe the logic is now correct for zero-length array sections.

[jdenny]

openmp/libomptarget/src/omptarget.cpp

492

Unfortunately, libomptarget does not have access to source code location. __tgt_ API functions do not have a related argument, like the ident_t *loc argument of __kmpc_ API functions in libomp. Of course, if there is demand for such a feature we can implement it, although we would need to change *EVERY* API function in libomptarget.

That would be nice, but in a later patch.

For this patch, perhaps just a simple message clarifying that the error is due to the present modifier would be sufficient. Should we define another macro next to FATAL_MESSAGE to print messages to stderr without exiting? The exit happens later when FATAL_MESSAGE is actually called.

[grokos]

openmp/libomptarget/src/omptarget.cpp

492

That would be good, such a message is missing from other places in libomptarget as well (e.g. when we explicitly try to extend an existing mapping). Currently such messages are only printed in debug output. It's important for the user to know why offload failed even if libomptarget is running in release mode.

[jdenny]

I added stderr messages for present violations, as discussed. I've named the new macro MESSAGE to contrast with FATAL_MESSAGE, but I'm open to suggestions.

I adjusted the new tests to check for these new messages instead of the LIBOMPTARGET_DEBUG=1 output.

[grokos]

The patch looks good now, thanks for all the work!

[jdenny]

Thanks for the review!

The Clang part of this (D83061) must land first, so I'll wait for that review.

[jdenny]

openmp/libomptarget/src/omptarget.cpp

490

This is right for omp target exit data, as in openmp/libomptarget/test/mapping/present/target_exit_data.c.

I think it's wrong for omp target data. My read of TR8 is that the present assertion does not happen on exit from a region. For example:

#pragma omp target data map(alloc:x)
{ 
  #pragma omp target data map(present,alloc:x)
  {
    #pragma omp target exit data map(delete:x)
  } // fails here but shouldn't
}

[jdenny]

openmp/libomptarget/src/omptarget.cpp

322

This comment is incorrect. We can also have PTR_AND_OBJ in the case of omp declare target link as discussed at http://lists.llvm.org/pipermail/openmp-dev/2020-July/003586.html.

[grokos]

openmp/libomptarget/src/omptarget.cpp

322

Right, the comment is not relevant anymore. But I think the conveyed message is still correct. Apart from structs/arrays, if we have a PTR_AND_OBJ for a global pointer then by definition the pointer itself has been mapped (since it's a global). Maybe update the comment? The logic is still correct I think.

[jdenny]

openmp/libomptarget/src/device.cpp

215

The \n in each of my MESSAGE calls should be dropped. MESSAGE already adds that.

[jdenny]

openmp/libomptarget/src/omptarget.cpp

322

Yes, I believe the code is right. I'll try to generalize the comment.

[grokos]

openmp/libomptarget/src/omptarget.cpp

322

I was just reminded that globals marked with declare target link are also PTR_AND_OBJ entries. declare target link already works in the exact same way as the proposed change we discussed at the July 15 telecon.

[jdenny]

Rebased onto latest D83061. Updated tests to use -fopenmp-version=51.

Removed extra newlines from MESSAGE calls.

Tried to fix PTR_AND_OBJ comment.

[jdenny]

openmp/libomptarget/src/omptarget.cpp

322

If we're so sure it's impossible to have !Pointer_TgtPtrBegin here, should we have an assert?

That way, if the situation evolves, testing will hopefully reveal when our assumptions here are wrong. Based on my understanding now, it would probably indicate a bug somewhere. But I guess it might just mean the assert then needs to be removed and the comment needs to be updated. Either way, it will be good to know about the change.

My understanding is that assert would only be active for debug builds. For release builds, we could still have the if and return OFFLOAD_FAIL for the sake of defensive programming.

490

In order to fix this, it looks like I could modify Clang to pass a new argument to __tgt_target_data_end to indicate whether it's just the end of a region (as in omp target exit data) or both the start and end of a region (as in omp target data).

Before I work on that, can somewhere confirm my understanding of TR8 behavior, described in the previous comment?

[jdenny]

Updated OMP_TGT_MAPTYPE_PRESENT to match new value of OMP_MAP_PRESENT in D83061.

[jdenny]

Rebased.

[grokos]

openmp/libomptarget/src/omptarget.cpp

322

Yes, we can include an assertion here.

329

It's declare target link, not plain declare target.

[jdenny]

openmp/libomptarget/src/omptarget.cpp

322

The assert would currently fail for the following:

int *x;
#pragma omp declare target link(x)

int main() {
  #pragma omp target map(present, tofrom: x[0:3])
  ;
  return 0;
}

[jdenny]

Tried to fix comments on PTR_AND_OBJ.

Added fixme for case of present at exit from omp target data.

openmp/libomptarget/src/omptarget.cpp

322

I didn't add the assert. I updated the comments to describe the situation as I understand it. Please advise if it's not right.

490

In order to fix this, it looks like I could modify Clang to pass a new argument to __tgt_target_data_end to indicate whether it's just the end of a region (as in omp target exit data) or both the start and end of a region (as in omp target data).

I reversed "omp target exit data" and "omp target data" there. Otherwise, I believe that comment is correct.

For now, I've decided to add a fixme and handle this in a later patch after pushing this series.

[grokos]

OK, let's proceed with this version now and later on we can fine-tune the behavior.

[jdenny]

Thanks again for the review.

[jdenny]

openmp/libomptarget/src/omptarget.cpp

491

It looks like DPxMOD and DPxPTR are not defined for release builds, and so bots are failing. I had only tested debug builds.

Here's an excerpt of the relevant code from omptarget.h:

#ifdef OMPTARGET_DEBUG
#include <stdio.h>
#define DEBUGP(prefix, ...)                                                    \
  {                                                                            \
    fprintf(stderr, "%s --> ", prefix);                                        \
    fprintf(stderr, __VA_ARGS__);                                              \
  }

#ifndef __STDC_FORMAT_MACROS
#define __STDC_FORMAT_MACROS
#endif

#include <inttypes.h>
#define DPxMOD "0x%0*" PRIxPTR
#define DPxPTR(ptr) ((int)(2*sizeof(uintptr_t))), ((uintptr_t) (ptr))

Any objection to moving relevant parts before the #ifdef OMPTARGET_DEBUG? Or is there a better solution?

[grokos]

Sure, go ahead!

openmp/libomptarget/include/omptarget.h

enum tgt_map_type {
   // private variable - not mapped
   OMP_TGT_MAPTYPE_PRIVATE         = 0x080,
   // copy by value - not mapped
   OMP_TGT_MAPTYPE_LITERAL         = 0x100,
   // mapping is implicit
   OMP_TGT_MAPTYPE_IMPLICIT        = 0x200,
   // copy data to device
   OMP_TGT_MAPTYPE_CLOSE           = 0x400,
+  // runtime error if not already allocated
+  OMP_TGT_MAPTYPE_PRESENT         = 0x1000,
   // member of struct, member given by [16 MSBs] - 1
   OMP_TGT_MAPTYPE_MEMBER_OF       = 0xffff000000000000
 };
 
 enum OpenMPOffloadingDeclareTargetFlags {
   /// Mark the entry as having a 'link' attribute.
   OMP_DECLARE_TARGET_LINK = 0x01,
   /// Mark the entry as being a global constructor.

openmp/libomptarget/src/device.h

struct DeviceTy {
   }
 
   // Return true if data can be copied to DstDevice directly
   bool isDataExchangable(const DeviceTy& DstDevice);
 
   uint64_t getMapEntryRefCnt(void *HstPtrBegin);
   LookupResult lookupMapping(void *HstPtrBegin, int64_t Size);
   void *getOrAllocTgtPtr(void *HstPtrBegin, void *HstPtrBase, int64_t Size,
-      bool &IsNew, bool &IsHostPtr, bool IsImplicit, bool UpdateRefCount = true,
-      bool HasCloseModifier = false);
+                         bool &IsNew, bool &IsHostPtr, bool IsImplicit,
+                         bool UpdateRefCount, bool HasCloseModifier,
+                         bool HasPresentModifier);
   void *getTgtPtrBegin(void *HstPtrBegin, int64_t Size);
   void *getTgtPtrBegin(void *HstPtrBegin, int64_t Size, bool &IsLast,
       bool UpdateRefCount, bool &IsHostPtr);
   int deallocTgtPtr(void *TgtPtrBegin, int64_t Size, bool ForceDelete,
                     bool HasCloseModifier = false);
   int associatePtr(void *HstPtrBegin, void *TgtPtrBegin, int64_t Size);
   int disassociatePtr(void *HstPtrBegin);
 
   // calls to RTL

openmp/libomptarget/src/device.cpp

 
 // Used by target_data_begin
 // Return the target pointer begin (where the data will be moved).
 // Allocate memory if this is the first occurrence of this mapping.
 // Increment the reference counter.
 // If NULL is returned, then either data allocation failed or the user tried
 // to do an illegal mapping.
 void *DeviceTy::getOrAllocTgtPtr(void *HstPtrBegin, void *HstPtrBase,
-    int64_t Size, bool &IsNew, bool &IsHostPtr, bool IsImplicit,
-    bool UpdateRefCount, bool HasCloseModifier) {
+                                 int64_t Size, bool &IsNew, bool &IsHostPtr,
+                                 bool IsImplicit, bool UpdateRefCount,
+                                 bool HasCloseModifier,
+                                 bool HasPresentModifier) {
   void *rc = NULL;
   IsHostPtr = false;
   IsNew = false;
   DataMapMtx.lock();
   LookupResult lr = lookupMapping(HstPtrBegin, Size);
 
   // Check if the pointer is contained.
   // If a variable is mapped to the device manually by the user - which would
DP("Mapping exists%s with HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD ", "
         "Size=%ld,%s RefCount=%s\n", (IsImplicit ? " (implicit)" : ""),
         DPxPTR(HstPtrBegin), DPxPTR(tp), Size,
         (UpdateRefCount ? " updated" : ""),
         HT.isRefCountInf() ? "INF" : std::to_string(HT.getRefCount()).c_str());
     rc = (void *)tp;
   } else if ((lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter) && !IsImplicit) {
     // Explicit extension of mapped data - not allowed.
     DP("Explicit extension of mapping is not allowed.\n");
-  } else if (Size) {
-    // 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.
-    if (RTLs->RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY &&
+  } else if (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.
+    if (Size) {
       DP("Return HstPtrBegin " DPxMOD " Size=%ld RefCount=%s\n",
-         DPxPTR((uintptr_t)HstPtrBegin), Size, (UpdateRefCount ? " updated" : ""));
+         DPxPTR((uintptr_t)HstPtrBegin), Size,
+         (UpdateRefCount ? " updated" : ""));
       IsHostPtr = true;
       rc = HstPtrBegin;
-    } else {
-      // If it is not contained and Size > 0 we should create a new entry for it.
-      IsNew = true;
-      uintptr_t tp = (uintptr_t)RTL->data_alloc(RTLDeviceID, Size, HstPtrBegin);
-      DP("Creating new map entry: HstBase=" DPxMOD ", HstBegin=" DPxMOD ", "
-         "HstEnd=" DPxMOD ", TgtBegin=" DPxMOD "\n", DPxPTR(HstPtrBase),
-         DPxPTR(HstPtrBegin), DPxPTR((uintptr_t)HstPtrBegin + Size), DPxPTR(tp));
-      HostDataToTargetMap.emplace(
-          HostDataToTargetTy((uintptr_t)HstPtrBase, (uintptr_t)HstPtrBegin,
-                             (uintptr_t)HstPtrBegin + Size, tp));
-      rc = (void *)tp;
-    }
-  }
+    }
+  } else if (HasPresentModifier) {
+    DP("Mapping required by 'present' map type modifier does not exist for "
+       "HstPtrBegin=" DPxMOD ", Size=%ld\n",
+       DPxPTR(HstPtrBegin), Size);
+    MESSAGE("device mapping required by 'present' map type modifier does not "
+            "exist for host address " DPxMOD " (%ld bytes)",
+            DPxPTR(HstPtrBegin), Size);
+  } else if (Size) {
+    // If it is not contained and Size > 0, we should create a new entry for it.
+       "HstEnd=" DPxMOD ", TgtBegin=" DPxMOD "\n",
+       DPxPTR(HstPtrBase), DPxPTR(HstPtrBegin),
+       DPxPTR((uintptr_t)HstPtrBegin + Size), DPxPTR(tp));
 
   DataMapMtx.unlock();
   return rc;
 }
 
 // Used by target_data_begin, target_data_end, target_data_update and target.
 // Return the target pointer begin (where the data will be moved).
 // Decrement the reference counter if called from target_data_end.

openmp/libomptarget/src/omptarget.cpp

for (int32_t i = 0; i < arg_num; ++i) {
     // Address of pointer on the host and device, respectively.
     void *Pointer_HstPtrBegin, *Pointer_TgtPtrBegin;
     bool IsNew, Pointer_IsNew;
     bool IsHostPtr = false;
     bool IsImplicit = arg_types[i] & OMP_TGT_MAPTYPE_IMPLICIT;
     // Force the creation of a device side copy of the data when:
     // a close map modifier was associated with a map that contained a to.
     bool HasCloseModifier = arg_types[i] & OMP_TGT_MAPTYPE_CLOSE;
+    bool HasPresentModifier = arg_types[i] & OMP_TGT_MAPTYPE_PRESENT;
     // UpdateRef is based on MEMBER_OF instead of TARGET_PARAM because if we
     // have reached this point via __tgt_target_data_begin and not __tgt_target
     // then no argument is marked as TARGET_PARAM ("omp target data map" is not
     // associated with a target region, so there are no target parameters). This
     // may be considered a hack, we could revise the scheme in the future.
     bool UpdateRef = !(arg_types[i] & OMP_TGT_MAPTYPE_MEMBER_OF);
     if (arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ) {
       DP("Has a pointer entry: \n");
-      // base is address of pointer.
-      Pointer_TgtPtrBegin = Device.getOrAllocTgtPtr(HstPtrBase, HstPtrBase,
-          sizeof(void *), Pointer_IsNew, IsHostPtr, IsImplicit, UpdateRef,
-          HasCloseModifier);
+      // Base is address of pointer.
+      //
+      // Usually, the pointer is already allocated by this time.  For example:
+      //
+      //   #pragma omp target map(s.p[0:N])
+      //
+      // 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 Pointer_TgtPtrBegin 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.
+      Pointer_TgtPtrBegin = Device.getOrAllocTgtPtr(
+          HstPtrBase, HstPtrBase, sizeof(void *), Pointer_IsNew, IsHostPtr,
+          IsImplicit, UpdateRef, HasCloseModifier, HasPresentModifier);
       if (!Pointer_TgtPtrBegin) {
-        DP("Call to getOrAllocTgtPtr returned null pointer (device failure or "
-            "illegal mapping).\n");
+        DP("Call to getOrAllocTgtPtr returned null pointer (%s).\n",
+           HasPresentModifier ? "'present' map type modifier"
+                              : "device failure or illegal mapping");
         return OFFLOAD_FAIL;
       }
       DP("There are %zu bytes allocated at target address " DPxMOD " - is%s new"
           "\n", sizeof(void *), DPxPTR(Pointer_TgtPtrBegin),
           (Pointer_IsNew ? "" : " not"));
       Pointer_HstPtrBegin = HstPtrBase;
       // modify current entry.
       HstPtrBase = *(void **)HstPtrBase;
       UpdateRef = true; // subsequently update ref count of pointee
     }
 
-    void *TgtPtrBegin = Device.getOrAllocTgtPtr(HstPtrBegin, HstPtrBase,
-        data_size, IsNew, IsHostPtr, IsImplicit, UpdateRef, HasCloseModifier);
-    if (!TgtPtrBegin && data_size) {
-      // If data_size==0, then the argument could be a zero-length pointer to
-      // NULL, so getOrAlloc() returning NULL is not an error.
-      DP("Call to getOrAllocTgtPtr returned null pointer (device failure or "
-          "illegal mapping).\n");
+    void *TgtPtrBegin = Device.getOrAllocTgtPtr(
+        HstPtrBegin, HstPtrBase, data_size, IsNew, IsHostPtr, IsImplicit,
+        UpdateRef, HasCloseModifier, HasPresentModifier);
+    if (!TgtPtrBegin && (data_size || HasPresentModifier)) {
+      DP("Call to getOrAllocTgtPtr returned null pointer (%s).\n",
+         HasPresentModifier ? "'present' map type modifier"
+                            : "device failure or illegal mapping");
       return OFFLOAD_FAIL;
     }
     DP("There are %" PRId64 " bytes allocated at target address " DPxMOD
         " - is%s new\n", data_size, DPxPTR(TgtPtrBegin),
         (IsNew ? "" : " not"));
 
     if (arg_types[i] & OMP_TGT_MAPTYPE_RETURN_PARAM) {
       uintptr_t Delta = (uintptr_t)HstPtrBegin - (uintptr_t)HstPtrBase;
if (member_of(arg_types[i]) < 0 && next_i < arg_num &&
       }
     }
 
     bool IsLast, IsHostPtr;
     bool UpdateRef = !(arg_types[i] & OMP_TGT_MAPTYPE_MEMBER_OF) ||
         (arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ);
     bool ForceDelete = arg_types[i] & OMP_TGT_MAPTYPE_DELETE;
     bool HasCloseModifier = arg_types[i] & OMP_TGT_MAPTYPE_CLOSE;
+    bool HasPresentModifier = arg_types[i] & OMP_TGT_MAPTYPE_PRESENT;
 
     // If PTR_AND_OBJ, HstPtrBegin is address of pointee
     void *TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBegin, data_size, IsLast,
         UpdateRef, IsHostPtr);
+    if (!TgtPtrBegin && (data_size || HasPresentModifier)) {
+      DP("Mapping does not exist (%s)\n",
+         (HasPresentModifier ? "'present' map type modifier" : "ignored"));
+      if (HasPresentModifier) {
+        // FIXME: This should not be an error on exit from "omp target data",
+        // but it should be an error upon entering an "omp target exit data".
+        MESSAGE("device mapping required by 'present' map type modifier does "
+                "not exist for host address " DPxMOD " (%ld bytes)",
+                DPxPTR(HstPtrBegin), data_size);
+        return OFFLOAD_FAIL;
+      }
+    } else {
       DP("There are %" PRId64 " bytes allocated at target address " DPxMOD
-        " - is%s last\n", data_size, DPxPTR(TgtPtrBegin),
-        (IsLast ? "" : " not"));
+         " - is%s last\n",
+         data_size, DPxPTR(TgtPtrBegin), (IsLast ? "" : " not"));
+    }
 
     bool DelEntry = IsLast || ForceDelete;
 
     if ((arg_types[i] & OMP_TGT_MAPTYPE_MEMBER_OF) &&
         !(arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ)) {
       DelEntry = false; // protect parent struct from being deallocated
     }
 

openmp/libomptarget/src/private.h

 typedef void (*MapperFuncPtrTy)(void *, void *, void *, int64_t, int64_t);
 
 // Function pointer type for target_data_* functions (target_data_begin,
 // target_data_end and target_data_update).
 typedef int (*TargetDataFuncPtrTy)(DeviceTy &, int32_t, void **, void **,
     int64_t *, int64_t *, void **, __tgt_async_info *);
 
 ////////////////////////////////////////////////////////////////////////////////
-// implementation for fatal messages
+// implementation for messages
 ////////////////////////////////////////////////////////////////////////////////
 
+#define MESSAGE0(_str)                                                         \
+  do {                                                                         \
+    fprintf(stderr, "Libomptarget message: %s\n", _str);                       \
+  } while (0)
+
+#define MESSAGE(_str, ...)                                                     \
+  do {                                                                         \
+    fprintf(stderr, "Libomptarget message: " _str "\n", __VA_ARGS__);          \
+  } while (0)
+
 #define FATAL_MESSAGE0(_num, _str)                                    \
   do {                                                                \
     fprintf(stderr, "Libomptarget fatal error %d: %s\n", _num, _str); \
     exit(1);                                                          \
   } while (0)
 
 #define FATAL_MESSAGE(_num, _str, ...)                              \
   do {                                                              \

openmp/libomptarget/test/mapping/present/target.c

+// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-fail-aarch64-unknown-linux-gnu 2>&1 \
+// RUN: | %fcheck-aarch64-unknown-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-fail-powerpc64-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-fail-powerpc64le-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64le-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-x86_64-pc-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-fail-x86_64-pc-linux-gnu 2>&1 \
+// RUN: | %fcheck-x86_64-pc-linux-gnu
+
+#include <stdio.h>
+
+int main() {
+  int i;
+
+  // CHECK: addr=0x[[#%x,HOST_ADDR:]], size=[[#%u,SIZE:]]
+  fprintf(stderr, "addr=%p, size=%ld\n", &i, sizeof i);
+
+  // CHECK-NOT: Libomptarget
+#pragma omp target data map(alloc: i)
+#pragma omp target map(present, alloc: i)
+  ;
+
+  // CHECK: i is present
+  fprintf(stderr, "i is present\n");
+
+  // 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 fatal error 1: failure of target construct while offloading is mandatory
+#pragma omp target map(present, alloc: i)
+  ;
+
+  // CHECK-NOT: i is present
+  fprintf(stderr, "i is present\n");
+
+  return 0;
+}

openmp/libomptarget/test/mapping/present/target_data.c

+// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-fail-aarch64-unknown-linux-gnu 2>&1 \
+// RUN: | %fcheck-aarch64-unknown-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-fail-powerpc64-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-fail-powerpc64le-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64le-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-x86_64-pc-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-fail-x86_64-pc-linux-gnu 2>&1 \
+// RUN: | %fcheck-x86_64-pc-linux-gnu
+
+#include <stdio.h>
+
+int main() {
+  int i;
+
+  // CHECK: addr=0x[[#%x,HOST_ADDR:]], size=[[#%u,SIZE:]]
+  fprintf(stderr, "addr=%p, size=%ld\n", &i, sizeof i);
+
+  // CHECK-NOT: Libomptarget
+#pragma omp target data map(alloc: i)
+#pragma omp target data map(present, alloc: i)
+  ;
+
+  // CHECK: i is present
+  fprintf(stderr, "i is present\n");
+
+  // 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 fatal error 1: failure of target construct while offloading is mandatory
+#pragma omp target data map(present, alloc: i)
+  ;
+
+  // CHECK-NOT: i is present
+  fprintf(stderr, "i is present\n");
+
+  return 0;
+}

openmp/libomptarget/test/mapping/present/target_enter_data.c

+// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-fail-aarch64-unknown-linux-gnu 2>&1 \
+// RUN: | %fcheck-aarch64-unknown-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-fail-powerpc64-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-fail-powerpc64le-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64le-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-x86_64-pc-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-fail-x86_64-pc-linux-gnu 2>&1 \
+// RUN: | %fcheck-x86_64-pc-linux-gnu
+
+#include <stdio.h>
+
+int main() {
+  int i;
+
+  // CHECK: addr=0x[[#%x,HOST_ADDR:]], size=[[#%u,SIZE:]]
+  fprintf(stderr, "addr=%p, size=%ld\n", &i, sizeof i);
+
+  // CHECK-NOT: Libomptarget
+#pragma omp target enter data map(alloc: i)
+#pragma omp target enter data map(present, alloc: i)
+#pragma omp target exit data map(delete: i)
+
+  // CHECK: i is present
+  fprintf(stderr, "i is present\n");
+
+  // 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 fatal error 1: failure of target construct while offloading is mandatory
+#pragma omp target enter data map(present, alloc: i)
+
+  // CHECK-NOT: i is present
+  fprintf(stderr, "i is present\n");
+
+  return 0;
+}

openmp/libomptarget/test/mapping/present/target_exit_data.c

+// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-fail-aarch64-unknown-linux-gnu 2>&1 \
+// RUN: | %fcheck-aarch64-unknown-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-fail-powerpc64-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-fail-powerpc64le-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64le-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-x86_64-pc-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-fail-x86_64-pc-linux-gnu 2>&1 \
+// RUN: | %fcheck-x86_64-pc-linux-gnu
+
+#include <stdio.h>
+
+int main() {
+  int i;
+
+  // CHECK: addr=0x[[#%x,HOST_ADDR:]], size=[[#%u,SIZE:]]
+  fprintf(stderr, "addr=%p, size=%ld\n", &i, sizeof i);
+
+  // CHECK-NOT: Libomptarget
+#pragma omp target enter data map(alloc: i)
+#pragma omp target exit data map(present, release: i)
+
+  // CHECK: i is present
+  fprintf(stderr, "i is present\n");
+
+  // 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 fatal error 1: failure of target construct while offloading is mandatory
+#pragma omp target exit data map(present, release: i)
+
+  // CHECK-NOT: i is present
+  fprintf(stderr, "i is present\n");
+
+  return 0;
+}

openmp/libomptarget/test/mapping/present/unified_shared_memory.c

+// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-aarch64-unknown-linux-gnu 2>&1 \
+// RUN: | %fcheck-aarch64-unknown-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-powerpc64-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-powerpc64le-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64le-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-x86_64-pc-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-x86_64-pc-linux-gnu 2>&1 \
+// RUN: | %fcheck-x86_64-pc-linux-gnu
+
+#include <stdio.h>
+
+// The runtime considers unified shared memory to be always present.
+#pragma omp requires unified_shared_memory
+
+int main() {
+  int i;
+
+  // CHECK-NOT: Libomptarget
+#pragma omp target data map(alloc: i)
+#pragma omp target map(present, alloc: i)
+  ;
+
+  // CHECK: i is present
+  fprintf(stderr, "i is present\n");
+
+  // CHECK-NOT: Libomptarget
+#pragma omp target map(present, alloc: i)
+  ;
+
+  // CHECK: is present
+  fprintf(stderr, "i is present\n");
+
+  return 0;
+}

openmp/libomptarget/test/mapping/present/zero_length_array_section.c

+// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-fail-aarch64-unknown-linux-gnu 2>&1 \
+// RUN: | %fcheck-aarch64-unknown-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-fail-powerpc64-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-fail-powerpc64le-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64le-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-x86_64-pc-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-fail-x86_64-pc-linux-gnu 2>&1 \
+// RUN: | %fcheck-x86_64-pc-linux-gnu
+
+#include <stdio.h>
+
+int main() {
+  int arr[5];
+
+  // CHECK: addr=0x[[#%x,HOST_ADDR:]]
+  fprintf(stderr, "addr=%p\n", arr);
+
+  // CHECK-NOT: Libomptarget
+#pragma omp target data map(alloc: arr[0:5])
+#pragma omp target map(present, alloc: arr[0:0])
+  ;
+
+  // CHECK: arr is present
+  fprintf(stderr, "arr is present\n");
+
+  // arr[0:0] doesn't create an actual mapping in the first directive.
+  //
+  // CHECK: Libomptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] (0 bytes)
+  // CHECK: Libomptarget fatal error 1: failure of target construct while offloading is mandatory
+#pragma omp target data map(alloc: arr[0:0])
+#pragma omp target map(present, alloc: arr[0:0])
+  ;
+
+  // CHECK-NOT: arr is present
+  fprintf(stderr, "arr is present\n");
+
+  return 0;
+}

openmp/libomptarget/test/mapping/present/zero_length_array_section_exit.c

+// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-fail-aarch64-unknown-linux-gnu 2>&1 \
+// RUN: | %fcheck-aarch64-unknown-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-fail-powerpc64-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-fail-powerpc64le-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64le-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-x86_64-pc-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-fail-x86_64-pc-linux-gnu 2>&1 \
+// RUN: | %fcheck-x86_64-pc-linux-gnu
+
+#include <stdio.h>
+
+int main() {
+  int arr[5];
+
+  // CHECK: addr=0x[[#%x,HOST_ADDR:]]
+  fprintf(stderr, "addr=%p\n", arr);
+
+  // CHECK-NOT: Libomptarget
+#pragma omp target enter data map(alloc: arr[0:5])
+#pragma omp target exit data map(present, release: arr[0:0])
+
+  // CHECK: arr is present
+  fprintf(stderr, "arr is present\n");
+
+  // arr[0:0] doesn't create an actual mapping in the first directive.
+  //
+  // CHECK: Libomptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] (0 bytes)
+  // CHECK: Libomptarget fatal error 1: failure of target construct while offloading is mandatory
+#pragma omp target enter data map(alloc: arr[0:0])
+#pragma omp target exit data map(present, release: arr[0:0])
+
+  // CHECK-NOT: arr is present
+  fprintf(stderr, "arr is present\n");
+
+  return 0;
+}