[OpenMP] Always emit debug messages that indicate offloading failure

[jhuber6]

Summary

This patch changes the libomptarget runtime to always emit debug messages that occur before offloading failure. The goal is to provide users with information about why their application failed in the target region rather than a single failure message. This is only done in regions that precede offloading failure so this should not impact runtime performance. if the debug environment variable is set then the message is forwarded to the debug output as usual.

A new environment variable was added for future use but does nothing in this current patch. LIBOMPTARGET_INFO will be used to report runtime information to the user if requrested, such as grid size, SPMD usage, or data mapping. It will take an integer indicating the level of information verbosity and a value of 0 will disable it.

[jhuber6]

This is the first iteration. I just took the messages and passed them to the user without modifying them in any way. I also quickly added the environment variable to output extra information to the user in the case of failure but haven't added any support for it.

I didn't add any messages for the offload failures from the plugins because I wasn't sure if I should be modifying those because they don't inherit directly from the rest of the libomptarget source. There's useful information that tends to be more specific about what failed so I can add that in if it's okay.

[jdoerfert]

LGTM, nice first step. We continue from here.

[ye-luo]

Please document the flags in the patch summary.

[jhuber6]

Fixed formatting. Some of the libomptarget test files seem to just randomly fail between tests, which makes it kind of annoying to test for correctness. Hopefully nothing breaks in the background.

openmp/libomptarget/src/api.cpp

 EXTERN int omp_target_memcpy(void *dst, void *src, size_t length,
     size_t dst_offset, size_t src_offset, int dst_device, int src_device) {
   DP("Call to omp_target_memcpy, dst device %d, src device %d, "
       "dst addr " DPxMOD ", src addr " DPxMOD ", dst offset %zu, "
       "src offset %zu, length %zu\n", dst_device, src_device, DPxPTR(dst),
       DPxPTR(src), dst_offset, src_offset, length);
 
   if (!dst || !src || length <= 0) {
-    DP("Call to omp_target_memcpy with invalid arguments\n");
+    REPORT("Call to omp_target_memcpy with invalid arguments\n");
     return OFFLOAD_FAIL;
   }
 
   if (src_device != omp_get_initial_device() && !device_is_ready(src_device)) {
-      DP("omp_target_memcpy returns OFFLOAD_FAIL\n");
+      REPORT("omp_target_memcpy returns OFFLOAD_FAIL\n");
       return OFFLOAD_FAIL;
   }
 
   if (dst_device != omp_get_initial_device() && !device_is_ready(dst_device)) {
-      DP("omp_target_memcpy returns OFFLOAD_FAIL\n");
+      REPORT("omp_target_memcpy returns OFFLOAD_FAIL\n");
       return OFFLOAD_FAIL;
   }
 
   int rc = OFFLOAD_SUCCESS;
   void *srcAddr = (char *)src + src_offset;
   void *dstAddr = (char *)dst + dst_offset;
 
   if (src_device == omp_get_initial_device() &&
EXTERN int omp_target_memcpy_rect(void *dst, void *src, size_t element_size,
   if (!(dst || src)) {
     DP("Call to omp_target_memcpy_rect returns max supported dimensions %d\n",
         INT_MAX);
     return INT_MAX;
   }
 
   if (!dst || !src || element_size < 1 || num_dims < 1 || !volume ||
       !dst_offsets || !src_offsets || !dst_dimensions || !src_dimensions) {
-    DP("Call to omp_target_memcpy_rect with invalid arguments\n");
+    REPORT("Call to omp_target_memcpy_rect with invalid arguments\n");
     return OFFLOAD_FAIL;
   }
 
   int rc;
   if (num_dims == 1) {
     rc = omp_target_memcpy(dst, src, element_size * volume[0],
         element_size * dst_offsets[0], element_size * src_offsets[0],
         dst_device, src_device);

 
 EXTERN int omp_target_associate_ptr(void *host_ptr, void *device_ptr,
     size_t size, size_t device_offset, int device_num) {
   DP("Call to omp_target_associate_ptr with host_ptr " DPxMOD ", "
       "device_ptr " DPxMOD ", size %zu, device_offset %zu, device_num %d\n",
       DPxPTR(host_ptr), DPxPTR(device_ptr), size, device_offset, device_num);
 
   if (!host_ptr || !device_ptr || size <= 0) {
-    DP("Call to omp_target_associate_ptr with invalid arguments\n");
+    REPORT("Call to omp_target_associate_ptr with invalid arguments\n");
     return OFFLOAD_FAIL;
   }
 
   if (device_num == omp_get_initial_device()) {
-    DP("omp_target_associate_ptr: no association possible on the host\n");
+    REPORT("omp_target_associate_ptr: no association possible on the host\n");
     return OFFLOAD_FAIL;
   }
 
   if (!device_is_ready(device_num)) {
-    DP("omp_target_associate_ptr returns OFFLOAD_FAIL\n");
+    REPORT("omp_target_associate_ptr returns OFFLOAD_FAIL\n");
     return OFFLOAD_FAIL;
   }
 
   DeviceTy& Device = Devices[device_num];
   void *device_addr = (void *)((uint64_t)device_ptr + (uint64_t)device_offset);
   int rc = Device.associatePtr(host_ptr, device_addr, size);
   DP("omp_target_associate_ptr returns %d\n", rc);
   return rc;
 }
 
 EXTERN int omp_target_disassociate_ptr(void *host_ptr, int device_num) {
   DP("Call to omp_target_disassociate_ptr with host_ptr " DPxMOD ", "
       "device_num %d\n", DPxPTR(host_ptr), device_num);
 
   if (!host_ptr) {
-    DP("Call to omp_target_associate_ptr with invalid host_ptr\n");
+    REPORT("Call to omp_target_associate_ptr with invalid host_ptr\n");
     return OFFLOAD_FAIL;
   }
 
   if (device_num == omp_get_initial_device()) {
-    DP("omp_target_disassociate_ptr: no association possible on the host\n");
+    REPORT("omp_target_disassociate_ptr: no association possible on the host\n");
     return OFFLOAD_FAIL;
   }
 
   if (!device_is_ready(device_num)) {
-    DP("omp_target_disassociate_ptr returns OFFLOAD_FAIL\n");
+    REPORT("omp_target_disassociate_ptr returns OFFLOAD_FAIL\n");
     return OFFLOAD_FAIL;
   }
 
   DeviceTy& Device = Devices[device_num];
   int rc = Device.disassociatePtr(host_ptr);
   DP("omp_target_disassociate_ptr returns %d\n", rc);
   return rc;
 }

openmp/libomptarget/src/device.cpp

if (search != HostDataToTargetMap.end()) {
     bool isValid = search->HstPtrEnd == (uintptr_t)HstPtrBegin + Size &&
                    search->TgtPtrBegin == (uintptr_t)TgtPtrBegin;
     DataMapMtx.unlock();
     if (isValid) {
       DP("Attempt to re-associate the same device ptr+offset with the same "
          "host ptr, nothing to do\n");
       return OFFLOAD_SUCCESS;
     } else {
-      DP("Not allowed to re-associate a different device ptr+offset with the "
-         "same host ptr\n");
+      REPORT("Not allowed to re-associate a different device ptr+offset with "
+             "the same host ptr\n");
       return OFFLOAD_FAIL;
     }
   }
 
   // Mapping does not exist, allocate it with refCount=INF
   HostDataToTargetTy newEntry((uintptr_t) HstPtrBegin /*HstPtrBase*/,
                               (uintptr_t) HstPtrBegin /*HstPtrBegin*/,
                               (uintptr_t) HstPtrBegin + Size /*HstPtrEnd*/,
int DeviceTy::disassociatePtr(void *HstPtrBegin) {
   if (search != HostDataToTargetMap.end()) {
     // Mapping exists
     if (search->isRefCountInf()) {
       DP("Association found, removing it\n");
       HostDataToTargetMap.erase(search);
       DataMapMtx.unlock();
       return OFFLOAD_SUCCESS;
     } else {
-      DP("Trying to disassociate a pointer which was not mapped via "
+      REPORT("Trying to disassociate a pointer which was not mapped via "
              "omp_target_associate_ptr\n");
     }
   }
 
   // Mapping not found
   DataMapMtx.unlock();
-  DP("Association not found\n");
+  REPORT("Association not found\n");
   return OFFLOAD_FAIL;
 }
 
 // Get ref count of map entry containing HstPtrBegin
 uint64_t DeviceTy::getMapEntryRefCnt(void *HstPtrBegin) {
   uintptr_t hp = (uintptr_t)HstPtrBegin;
   uint64_t RefCnt = 0;
 
if (HT.decRefCount() == 0) {
       deleteData((void *)HT.TgtPtrBegin);
       DP("Removing%s mapping with HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD
           ", Size=%" PRId64 "\n", (ForceDelete ? " (forced)" : ""),
           DPxPTR(HT.HstPtrBegin), DPxPTR(HT.TgtPtrBegin), Size);
       HostDataToTargetMap.erase(lr.Entry);
     }
     rc = OFFLOAD_SUCCESS;
   } else {
-    DP("Section to delete (hst addr " DPxMOD ") does not exist in the allocated"
-       " memory\n", DPxPTR(HstPtrBegin));
+    REPORT("Section to delete (hst addr " DPxMOD ") does not exist in the"
+           " allocated memory\n", DPxPTR(HstPtrBegin));
     rc = OFFLOAD_FAIL;
   }
 
   DataMapMtx.unlock();
   return rc;
 }
 
 /// Init device, should not be called directly.

openmp/libomptarget/src/interface.cpp

case tgt_disabled:
       }
       break;
     case tgt_default:
       FATAL_MESSAGE0(1, "default offloading policy must be switched to "
                         "mandatory or disabled");
       break;
     case tgt_mandatory:
       if (!success) {
+        if (InfoLevel > 0)
+          MESSAGE0("LIBOMPTARGET_INFO is not supported yet");
         FATAL_MESSAGE0(1, "failure of target construct while offloading is mandatory");
       }
       break;
   }
 }
 
 ////////////////////////////////////////////////////////////////////////////////
 /// adds requires flags
EXTERN int __tgt_target_mapper(int64_t device_id, void *host_ptr,
   DP("Entering target region with entry point " DPxMOD " and device Id %"
       PRId64 "\n", DPxPTR(host_ptr), device_id);
 
   if (device_id == OFFLOAD_DEVICE_DEFAULT) {
     device_id = omp_get_default_device();
   }
 
   if (CheckDeviceAndCtors(device_id) != OFFLOAD_SUCCESS) {
-    DP("Failed to get device %" PRId64 " ready\n", device_id);
+    REPORT("Failed to get device %" PRId64 " ready\n", device_id);
     HandleTargetOutcome(false);
     return OFFLOAD_FAIL;
   }
 
 #ifdef OMPTARGET_DEBUG
   for (int i=0; i<arg_num; ++i) {
     DP("Entry %2d: Base=" DPxMOD ", Begin=" DPxMOD ", Size=%" PRId64
         ", Type=0x%" PRIx64 "\n", i, DPxPTR(args_base[i]), DPxPTR(args[i]),
EXTERN int __tgt_target_teams_mapper(int64_t device_id, void *host_ptr,
   DP("Entering target region with entry point " DPxMOD " and device Id %"
       PRId64 "\n", DPxPTR(host_ptr), device_id);
 
   if (device_id == OFFLOAD_DEVICE_DEFAULT) {
     device_id = omp_get_default_device();
   }
 
   if (CheckDeviceAndCtors(device_id) != OFFLOAD_SUCCESS) {
-    DP("Failed to get device %" PRId64 " ready\n", device_id);
+    REPORT("Failed to get device %" PRId64 " ready\n", device_id);
            HandleTargetOutcome(false);
     return OFFLOAD_FAIL;
   }
 
 #ifdef OMPTARGET_DEBUG
   for (int i=0; i<arg_num; ++i) {
     DP("Entry %2d: Base=" DPxMOD ", Begin=" DPxMOD ", Size=%" PRId64
         ", Type=0x%" PRIx64 "\n", i, DPxPTR(args_base[i]), DPxPTR(args[i]),
         arg_sizes[i], arg_types[i]);

openmp/libomptarget/src/omptarget.cpp

 #include "rtl.h"
 
 #include <cassert>
 #include <vector>
 
 #ifdef OMPTARGET_DEBUG
 int DebugLevel = 0;
 #endif // OMPTARGET_DEBUG
-
+int InfoLevel = 0;
 
 
 /* All begin addresses for partially mapped structs must be 8-aligned in order
  * to ensure proper alignment of members. E.g.
  *
  * struct S {
  *   int a;   // 4-aligned
  *   int b;   // 4-aligned
  *   int *p;  // 8-aligned
  * } s1;
if (TransTable->TargetsTable[device_id] != 0) {
       continue;
     }
 
     // 1) get image.
     assert(TransTable->TargetsImages.size() > (size_t)device_id &&
            "Not expecting a device ID outside the table's bounds!");
     __tgt_device_image *img = TransTable->TargetsImages[device_id];
     if (!img) {
-      DP("No image loaded for device id %d.\n", device_id);
+      REPORT("No image loaded for device id %d.\n", device_id);
       rc = OFFLOAD_FAIL;
       break;
     }
     // 2) load image into the target table.
     __tgt_target_table *TargetTable =
         TransTable->TargetsTable[device_id] = Device.load_binary(img);
     // Unable to get table for this image: invalidate image and fail.
     if (!TargetTable) {
-      DP("Unable to generate entries table for device id %d.\n", device_id);
+      REPORT("Unable to generate entries table for device id %d.\n", device_id);
       TransTable->TargetsImages[device_id] = 0;
       rc = OFFLOAD_FAIL;
       break;
     }
 
     // Verify whether the two table sizes match.
     size_t hsize =
         TransTable->HostTable.EntriesEnd - TransTable->HostTable.EntriesBegin;
     size_t tsize = TargetTable->EntriesEnd - TargetTable->EntriesBegin;
 
     // Invalid image for these host entries!
     if (hsize != tsize) {
-      DP("Host and Target tables mismatch for device id %d [%zx != %zx].\n",
+      REPORT("Host and Target tables mismatch for device id %d [%zx != %zx].\n",
          device_id, hsize, tsize);
       TransTable->TargetsImages[device_id] = 0;
       TransTable->TargetsTable[device_id] = 0;
       rc = OFFLOAD_FAIL;
       break;
     }
 
     // process global data that needs to be mapped.
     Device.DataMapMtx.lock();
if (!Device.PendingCtorsDtors.empty()) {
     for (auto &lib : Device.PendingCtorsDtors) {
       if (!lib.second.PendingCtors.empty()) {
         DP("Has pending ctors... call now\n");
         for (auto &entry : lib.second.PendingCtors) {
           void *ctor = entry;
           int rc = target(device_id, ctor, 0, NULL, NULL, NULL, NULL, NULL, 1,
               1, true /*team*/);
           if (rc != OFFLOAD_SUCCESS) {
-            DP("Running ctor " DPxMOD " failed.\n", DPxPTR(ctor));
+            REPORT("Running ctor " DPxMOD " failed.\n", DPxPTR(ctor));
             Device.PendingGlobalsMtx.unlock();
             return OFFLOAD_FAIL;
           }
         }
         // Clear the list to indicate that this device has been used
         lib.second.PendingCtors.clear();
         DP("Done with pending ctors for lib " DPxMOD "\n", DPxPTR(lib.first));
       }
     }
   }
   Device.HasPendingGlobals = false;
   Device.PendingGlobalsMtx.unlock();
 
   return OFFLOAD_SUCCESS;
 }
 
 // Check whether a device has been initialized, global ctors have been
 // executed and global data has been mapped; do so if not already done.
 int CheckDeviceAndCtors(int64_t device_id) {
   // Is device ready?
   if (!device_is_ready(device_id)) {
-    DP("Device %" PRId64 " is not ready.\n", device_id);
+    REPORT("Device %" PRId64 " is not ready.\n", device_id);
     return OFFLOAD_FAIL;
   }
 
   // Get device info.
   DeviceTy &Device = Devices[device_id];
 
   // Check whether global data has been mapped for this device
   Device.PendingGlobalsMtx.lock();
   bool hasPendingGlobals = Device.HasPendingGlobals;
   Device.PendingGlobalsMtx.unlock();
   if (hasPendingGlobals && InitLibrary(Device) != OFFLOAD_SUCCESS) {
-    DP("Failed to init globals on device %" PRId64 "\n", device_id);
+    REPORT("Failed to init globals on device %" PRId64 "\n", device_id);
     return OFFLOAD_FAIL;
   }
 
   return OFFLOAD_SUCCESS;
 }
 
 static int32_t getParentIndex(int64_t type) {
   return ((type & OMP_TGT_MAPTYPE_MEMBER_OF) >> 48) - 1;
if (arg_mappers && arg_mappers[i]) {
       // targetDataMapper variant which will call targetDataBegin again
       // with new arguments.
       DP("Calling targetDataMapper for the %dth argument\n", i);
 
       int rc = targetDataMapper(Device, args_base[i], args[i], arg_sizes[i],
                                 arg_types[i], arg_mappers[i], targetDataBegin);
 
       if (rc != OFFLOAD_SUCCESS) {
-        DP("Call to targetDataBegin via targetDataMapper for custom mapper"
+        REPORT("Call to targetDataBegin via targetDataMapper for custom mapper"
                " failed.\n");
         return OFFLOAD_FAIL;
       }
 
       // Skip the rest of this function, continue to the next argument.
       continue;
     }
 
     void *HstPtrBegin = args[i];
if (arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ) {
       // 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.
       PointerTgtPtrBegin = Device.getOrAllocTgtPtr(
           HstPtrBase, HstPtrBase, sizeof(void *), Pointer_IsNew, IsHostPtr,
           IsImplicit, UpdateRef, HasCloseModifier, HasPresentModifier);
       if (!PointerTgtPtrBegin) {
-        DP("Call to getOrAllocTgtPtr returned null pointer (%s).\n",
+        REPORT("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(PointerTgtPtrBegin),
           (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, HasPresentModifier);
     // If data_size==0, then the argument could be a zero-length pointer to
     // NULL, so getOrAlloc() returning NULL is not an error.
     if (!TgtPtrBegin && (data_size || HasPresentModifier)) {
-      DP("Call to getOrAllocTgtPtr returned null pointer (%s).\n",
+      REPORT("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 (arg_types[i] & OMP_TGT_MAPTYPE_TO) {
       }
 
       if (copy && !IsHostPtr) {
         DP("Moving %" PRId64 " bytes (hst:" DPxMOD ") -> (tgt:" DPxMOD ")\n",
            data_size, DPxPTR(HstPtrBegin), DPxPTR(TgtPtrBegin));
         int rt = Device.submitData(TgtPtrBegin, HstPtrBegin, data_size,
                                    async_info_ptr);
         if (rt != OFFLOAD_SUCCESS) {
-          DP("Copying data to device failed.\n");
+          REPORT("Copying data to device failed.\n");
           return OFFLOAD_FAIL;
         }
       }
     }
 
     if (arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ && !IsHostPtr) {
       DP("Update pointer (" DPxMOD ") -> [" DPxMOD "]\n",
          DPxPTR(PointerTgtPtrBegin), DPxPTR(TgtPtrBegin));
       uint64_t Delta = (uint64_t)HstPtrBegin - (uint64_t)HstPtrBase;
       void *TgtPtrBase = (void *)((uint64_t)TgtPtrBegin - Delta);
       int rt = Device.submitData(PointerTgtPtrBegin, &TgtPtrBase,
                                  sizeof(void *), async_info_ptr);
       if (rt != OFFLOAD_SUCCESS) {
-        DP("Copying data to device failed.\n");
+        REPORT("Copying data to device failed.\n");
         return OFFLOAD_FAIL;
       }
       // create shadow pointers for this entry
       Device.ShadowMtx.lock();
       Device.ShadowPtrMap[Pointer_HstPtrBegin] = {
           HstPtrBase, PointerTgtPtrBegin, TgtPtrBase};
       Device.ShadowMtx.unlock();
     }
if (ArgMappers && ArgMappers[I]) {
       // targetDataMapper variant which will call targetDataEnd again
       // with new arguments.
       DP("Calling targetDataMapper for the %dth argument\n", I);
 
       Ret = targetDataMapper(Device, ArgBases[I], Args[I], ArgSizes[I],
                              ArgTypes[I], ArgMappers[I], targetDataEnd);
 
       if (Ret != OFFLOAD_SUCCESS) {
-        DP("Call to targetDataEnd via targetDataMapper for custom mapper"
+        REPORT("Call to targetDataEnd via targetDataMapper for custom mapper"
                " failed.\n");
         return OFFLOAD_FAIL;
       }
 
       // Skip the rest of this function, continue to the next argument.
       continue;
     }
 
     void *HstPtrBegin = Args[I];
if ((ArgTypes[I] & OMP_TGT_MAPTYPE_FROM) || DelEntry) {
         if ((DelEntry || Always || CopyMember) &&
             !(RTLs->RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY &&
               TgtPtrBegin == HstPtrBegin)) {
           DP("Moving %" PRId64 " bytes (tgt:" DPxMOD ") -> (hst:" DPxMOD ")\n",
              DataSize, DPxPTR(TgtPtrBegin), DPxPTR(HstPtrBegin));
           Ret = Device.retrieveData(HstPtrBegin, TgtPtrBegin, DataSize,
                                     AsyncInfo);
           if (Ret != OFFLOAD_SUCCESS) {
-            DP("Copying data from device failed.\n");
+            REPORT("Copying data from device failed.\n");
             return OFFLOAD_FAIL;
           }
         }
       }
 
       // If we copied back to the host a struct/array containing pointers, we
       // need to restore the original host pointer values from their shadow
       // copies. If the struct is going to be deallocated, remove any remaining
int targetDataEnd(DeviceTy &Device, int32_t ArgNum, void **ArgBases,
   // If AsyncInfo is nullptr, the previous data transfer (if has) will be
   // synchronous, so we don't need to synchronize again. If AsyncInfo->Queue is
   // nullptr, there is no data transfer happened because once there is,
   // AsyncInfo->Queue will not be nullptr, so again, we don't need to
   // synchronize.
   if (AsyncInfo && AsyncInfo->Queue) {
     Ret = Device.synchronize(AsyncInfo);
     if (Ret != OFFLOAD_SUCCESS) {
-      DP("Failed to synchronize device.\n");
+      REPORT("Failed to synchronize device.\n");
       return OFFLOAD_FAIL;
     }
   }
 
   // Deallocate target pointer
   for (DeallocTgtPtrInfo &Info : DeallocTgtPtrs) {
     Ret = Device.deallocTgtPtr(Info.HstPtrBegin, Info.DataSize,
                                Info.ForceDelete, Info.HasCloseModifier);
     if (Ret != OFFLOAD_SUCCESS) {
-      DP("Deallocating data from device failed.\n");
+      REPORT("Deallocating data from device failed.\n");
       return OFFLOAD_FAIL;
     }
   }
 
   return OFFLOAD_SUCCESS;
 }
 
 /// Internal function to pass data to/from the target.
if (arg_mappers && arg_mappers[i]) {
       // with new arguments.
       DP("Calling targetDataMapper for the %dth argument\n", i);
 
       int rc =
           targetDataMapper(Device, args_base[i], args[i], arg_sizes[i],
                            arg_types[i], arg_mappers[i], target_data_update);
 
       if (rc != OFFLOAD_SUCCESS) {
-        DP("Call to target_data_update via targetDataMapper for custom mapper"
+        REPORT("Call to target_data_update via targetDataMapper for custom mapper"
                " failed.\n");
         return OFFLOAD_FAIL;
       }
 
       // Skip the rest of this function, continue to the next argument.
       continue;
     }
 
     void *HstPtrBegin = args[i];
if (RTLs->RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY &&
       continue;
     }
 
     if (arg_types[i] & OMP_TGT_MAPTYPE_FROM) {
       DP("Moving %" PRId64 " bytes (tgt:" DPxMOD ") -> (hst:" DPxMOD ")\n",
           arg_sizes[i], DPxPTR(TgtPtrBegin), DPxPTR(HstPtrBegin));
       int rt = Device.retrieveData(HstPtrBegin, TgtPtrBegin, MapSize, nullptr);
       if (rt != OFFLOAD_SUCCESS) {
-        DP("Copying data from device failed.\n");
+        REPORT("Copying data from device failed.\n");
         return OFFLOAD_FAIL;
       }
 
       uintptr_t lb = (uintptr_t) HstPtrBegin;
       uintptr_t ub = (uintptr_t) HstPtrBegin + MapSize;
       Device.ShadowMtx.lock();
       for (ShadowPtrListTy::iterator it = Device.ShadowPtrMap.begin();
           it != Device.ShadowPtrMap.end(); ++it) {
if (arg_types[i] & OMP_TGT_MAPTYPE_FROM) {
       Device.ShadowMtx.unlock();
     }
 
     if (arg_types[i] & OMP_TGT_MAPTYPE_TO) {
       DP("Moving %" PRId64 " bytes (hst:" DPxMOD ") -> (tgt:" DPxMOD ")\n",
           arg_sizes[i], DPxPTR(HstPtrBegin), DPxPTR(TgtPtrBegin));
       int rt = Device.submitData(TgtPtrBegin, HstPtrBegin, MapSize, nullptr);
       if (rt != OFFLOAD_SUCCESS) {
-        DP("Copying data to device failed.\n");
+        REPORT("Copying data to device failed.\n");
         return OFFLOAD_FAIL;
       }
 
       uintptr_t lb = (uintptr_t) HstPtrBegin;
       uintptr_t ub = (uintptr_t) HstPtrBegin + MapSize;
       Device.ShadowMtx.lock();
       for (ShadowPtrListTy::iterator it = Device.ShadowPtrMap.begin();
           it != Device.ShadowPtrMap.end(); ++it) {
         void **ShadowHstPtrAddr = (void **)it->first;
         if ((uintptr_t)ShadowHstPtrAddr < lb)
           continue;
         if ((uintptr_t)ShadowHstPtrAddr >= ub)
           break;
         DP("Restoring original target pointer value " DPxMOD " for target "
            "pointer " DPxMOD "\n",
            DPxPTR(it->second.TgtPtrVal), DPxPTR(it->second.TgtPtrAddr));
         rt = Device.submitData(it->second.TgtPtrAddr, &it->second.TgtPtrVal,
                                sizeof(void *), nullptr);
         if (rt != OFFLOAD_SUCCESS) {
-          DP("Copying data to device failed.\n");
+          REPORT("Copying data to device failed.\n");
           Device.ShadowMtx.unlock();
           return OFFLOAD_FAIL;
         }
       }
       Device.ShadowMtx.unlock();
     }
   }
   return OFFLOAD_SUCCESS;
int processDataBefore(int64_t DeviceId, void *HostPtr, int32_t ArgNum,
                       std::vector<void *> &TgtArgs,
                       std::vector<ptrdiff_t> &TgtOffsets,
                       PrivateArgumentManagerTy &PrivateArgumentManager,
                       __tgt_async_info *AsyncInfo) {
   DeviceTy &Device = Devices[DeviceId];
   int Ret = targetDataBegin(Device, ArgNum, ArgBases, Args, ArgSizes, ArgTypes,
                             ArgMappers, AsyncInfo);
   if (Ret != OFFLOAD_SUCCESS) {
-    DP("Call to targetDataBegin failed, abort target.\n");
+    REPORT("Call to targetDataBegin failed, abort target.\n");
     return OFFLOAD_FAIL;
   }
 
   // List of (first-)private arrays allocated for this target region
   std::vector<int> TgtArgsPositions(ArgNum, -1);
 
   for (int32_t I = 0; I < ArgNum; ++I) {
     if (!(ArgTypes[I] & OMP_TGT_MAPTYPE_TARGET_PARAM)) {
if (!(ArgTypes[I] & OMP_TGT_MAPTYPE_TARGET_PARAM)) {
              DPxPTR(HstPtrVal));
           continue;
         }
         DP("Update lambda reference (" DPxMOD ") -> [" DPxMOD "]\n",
            DPxPTR(PointerTgtPtrBegin), DPxPTR(TgtPtrBegin));
         Ret = Device.submitData(TgtPtrBegin, &PointerTgtPtrBegin,
                                 sizeof(void *), AsyncInfo);
         if (Ret != OFFLOAD_SUCCESS) {
-          DP("Copying data to device failed.\n");
+          REPORT("Copying data to device failed.\n");
           return OFFLOAD_FAIL;
         }
       }
       continue;
     }
     void *HstPtrBegin = Args[I];
     void *HstPtrBase = ArgBases[I];
     void *TgtPtrBegin;
     ptrdiff_t TgtBaseOffset;
     bool IsLast, IsHostPtr; // unused.
     if (ArgTypes[I] & OMP_TGT_MAPTYPE_LITERAL) {
       DP("Forwarding first-private value " DPxMOD " to the target construct\n",
          DPxPTR(HstPtrBase));
       TgtPtrBegin = HstPtrBase;
       TgtBaseOffset = 0;
     } else if (ArgTypes[I] & OMP_TGT_MAPTYPE_PRIVATE) {
       TgtBaseOffset = (intptr_t)HstPtrBase - (intptr_t)HstPtrBegin;
       const bool IsFirstPrivate = ArgTypes[I] & OMP_TGT_MAPTYPE_TO;
       Ret = PrivateArgumentManager.addArg(HstPtrBegin, ArgSizes[I],
                                           TgtBaseOffset, IsFirstPrivate,
                                           TgtPtrBegin, TgtArgs.size());
       if (Ret != OFFLOAD_SUCCESS) {
-        DP("Failed to process %sprivate argument " DPxMOD "\n",
+        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);
       TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBegin, ArgSizes[I], IsLast,
                                           false, IsHostPtr);
       TgtBaseOffset = (intptr_t)HstPtrBase - (intptr_t)HstPtrBegin;
int processDataAfter(int64_t DeviceId, void *HostPtr, int32_t ArgNum,
                      PrivateArgumentManagerTy &PrivateArgumentManager,
                      __tgt_async_info *AsyncInfo) {
   DeviceTy &Device = Devices[DeviceId];
 
   // Move data from device.
   int Ret = targetDataEnd(Device, ArgNum, ArgBases, Args, ArgSizes, ArgTypes,
                           ArgMappers, AsyncInfo);
   if (Ret != OFFLOAD_SUCCESS) {
-    DP("Call to targetDataEnd failed, abort targe.\n");
+    REPORT("Call to targetDataEnd failed, abort target.\n");
     return OFFLOAD_FAIL;
   }
 
   // Free target memory for private arguments
   Ret = PrivateArgumentManager.free();
   if (Ret != OFFLOAD_SUCCESS) {
-    DP("Failed to deallocate target memory for private args\n");
+    REPORT("Failed to deallocate target memory for private args\n");
     return OFFLOAD_FAIL;
   }
 
   return OFFLOAD_SUCCESS;
 }
 } // namespace
 
 /// performs the same actions as data_begin in case arg_num is
 /// non-zero and initiates run of the offloaded region on the target platform;
 /// if arg_num is non-zero after the region execution is done it also
 /// performs the same action as data_update and data_end above. This function
 /// returns 0 if it was able to transfer the execution to a target and an
 /// integer different from zero otherwise.
 int target(int64_t DeviceId, void *HostPtr, int32_t ArgNum, void **ArgBases,
            void **Args, int64_t *ArgSizes, int64_t *ArgTypes, void **ArgMappers,
            int32_t TeamNum, int32_t ThreadLimit, int IsTeamConstruct) {
   DeviceTy &Device = Devices[DeviceId];
 
   TableMap *TM = getTableMap(HostPtr);
   // No map for this host pointer found!
   if (!TM) {
-    DP("Host ptr " DPxMOD " does not have a matching target pointer.\n",
+    REPORT("Host ptr " DPxMOD " does not have a matching target pointer.\n",
        DPxPTR(HostPtr));
     return OFFLOAD_FAIL;
   }
 
   // get target table.
   __tgt_target_table *TargetTable = nullptr;
   {
     std::lock_guard<std::mutex> TrlTblLock(*TrlTblMtx);
     assert(TM->Table->TargetsTable.size() > (size_t)DeviceId &&
int target(int64_t DeviceId, void *HostPtr, int32_t ArgNum, void **ArgBases,
 
   PrivateArgumentManagerTy PrivateArgumentManager(Device, &AsyncInfo);
 
   // Process data, such as data mapping, before launching the kernel
   int Ret = processDataBefore(DeviceId, HostPtr, ArgNum, ArgBases, Args,
                               ArgSizes, ArgTypes, ArgMappers, TgtArgs,
                               TgtOffsets, PrivateArgumentManager, &AsyncInfo);
   if (Ret != OFFLOAD_SUCCESS) {
-    DP("Failed to process data before launching the kernel.\n");
+    REPORT("Failed to process data before launching the kernel.\n");
     return OFFLOAD_FAIL;
   }
 
   // Get loop trip count
   uint64_t LoopTripCount = getLoopTripCount(DeviceId);
 
   // Launch device execution.
   void *TgtEntryPtr = TargetTable->EntriesBegin[TM->Index].addr;
   DP("Launching target execution %s with pointer " DPxMOD " (index=%d).\n",
      TargetTable->EntriesBegin[TM->Index].name, DPxPTR(TgtEntryPtr), TM->Index);
 
   if (IsTeamConstruct)
     Ret = Device.runTeamRegion(TgtEntryPtr, &TgtArgs[0], &TgtOffsets[0],
                                TgtArgs.size(), TeamNum, ThreadLimit,
                                LoopTripCount, &AsyncInfo);
   else
     Ret = Device.runRegion(TgtEntryPtr, &TgtArgs[0], &TgtOffsets[0],
                            TgtArgs.size(), &AsyncInfo);
 
   if (Ret != OFFLOAD_SUCCESS) {
-    DP("Executing target region abort target.\n");
+    REPORT("Executing target region abort target.\n");
     return OFFLOAD_FAIL;
   }
 
   // Transfer data back and deallocate target memory for (first-)private
   // variables
   Ret = processDataAfter(DeviceId, HostPtr, ArgNum, ArgBases, Args, ArgSizes,
                          ArgTypes, ArgMappers, PrivateArgumentManager,
                          &AsyncInfo);
   if (Ret != OFFLOAD_SUCCESS) {
-    DP("Failed to process data after launching the kernel.\n");
+    REPORT("Failed to process data after launching the kernel.\n");
     return OFFLOAD_FAIL;
   }
 
   return OFFLOAD_SUCCESS;
 }

openmp/libomptarget/src/private.h

 
 #ifndef _OMPTARGET_PRIVATE_H
 #define _OMPTARGET_PRIVATE_H
 
 #include <omptarget.h>
 
 #include <cstdint>
 
 extern int targetDataBegin(DeviceTy &Device, int32_t arg_num, void **args_base,
                            void **args, int64_t *arg_sizes, int64_t *arg_types,
                            void **arg_mappers,
                            __tgt_async_info *async_info_ptr);
 
 extern int targetDataEnd(DeviceTy &Device, int32_t ArgNum, void **ArgBases,
                          void **Args, int64_t *ArgSizes, int64_t *ArgTypes,
                          void **ArgMappers, __tgt_async_info *AsyncInfo);
 
 extern int target_data_update(DeviceTy &Device, int32_t arg_num,
                               void **args_base, void **args,
                               int64_t *arg_sizes, int64_t *arg_types,
                               void **arg_mappers,
                               __tgt_async_info *async_info_ptr = nullptr);
 
 extern int target(int64_t DeviceId, void *HostPtr, int32_t ArgNum,
                   void **ArgBases, void **Args, int64_t *ArgSizes,
                   int64_t *ArgTypes, void **ArgMappers, int32_t TeamNum,
struct MapComponentInfoTy {
   MapComponentInfoTy(void *Base, void *Begin, int64_t Size, int64_t Type)
       : Base(Base), Begin(Begin), Size(Size), Type(Type) {}
 };
 
 // This structure stores all components of a user-defined mapper. The number of
 // components are dynamically decided, so we utilize C++ STL vector
 // implementation here.
 struct MapperComponentsTy {
   std::vector<MapComponentInfoTy> Components;
   int32_t size() { return Components.size(); }
 };
 
 // The mapper function pointer type. It follows the signature below:
 // void .omp_mapper.<type_name>.<mapper_id>.(void *rt_mapper_handle,
 //                                           void *base, void *begin,
 //                                           size_t size, int64_t type);
 typedef void (*MapperFuncPtrTy)(void *, void *, void *, int64_t, int64_t);
 
 // Function pointer type for target_data_* functions (targetDataBegin,
 // targetDataEnd and target_data_update).
 typedef int (*TargetDataFuncPtrTy)(DeviceTy &, int32_t, void **, void **,
                                    int64_t *, int64_t *, void **,
                                    __tgt_async_info *);
 
 ////////////////////////////////////////////////////////////////////////////////
 // implementation for messages
 ////////////////////////////////////////////////////////////////////////////////
 
 #define MESSAGE0(_str)                                                         \

 
 #define FATAL_MESSAGE(_num, _str, ...)                                         \
   do {                                                                         \
     fprintf(stderr, "Libomptarget fatal error %d:" _str "\n", _num,            \
             __VA_ARGS__);                                                      \
     abort();                                                                   \
   } while (0)
 
+#define FAILURE_MESSAGE(...)                                                   \
+  do {                                                                         \
+    fprintf(stderr, "Libomptarget error: ");                                   \
+    fprintf(stderr, __VA_ARGS__);                                              \
+  } while (0)
+
 // Implemented in libomp, they are called from within __tgt_* functions.
 #ifdef __cplusplus
 extern "C" {
 #endif
 // functions that extract info from libomp; keep in sync
 int omp_get_default_device(void) __attribute__((weak));
 int32_t __kmpc_omp_taskwait(void *loc_ref, int32_t gtid) __attribute__((weak));
 int32_t __kmpc_global_thread_num(void *) __attribute__((weak));
 int __kmpc_get_target_offload(void) __attribute__((weak));
 #ifdef __cplusplus
 }
 #endif
 
+extern int InfoLevel;
 #ifdef OMPTARGET_DEBUG
 extern int DebugLevel;
 
 #define DP(...) \
   do { \
     if (DebugLevel > 0) { \
       DEBUGP("Libomptarget", __VA_ARGS__); \
     } \
   } while (false)
 #else // OMPTARGET_DEBUG
 #define DP(...) {}
 #endif // OMPTARGET_DEBUG
 
+// Report debug messages that result in offload failure always
+#ifdef OMPTARGET_DEBUG
+#define REPORT(...) \
+  do { \
+    if (DebugLevel > 0) { \
+      DP(__VA_ARGS__); \
+    } else { \
+      FAILURE_MESSAGE(__VA_ARGS__); \
+    } \
+  } while (false)
+#else
+#define REPORT(...) FAILURE_MESSAGE(__VA_ARGS__);
+#endif
+
 #endif

openmp/libomptarget/src/rtl.cpp

__attribute__((destructor(101))) void deinit() {
   delete RTLsMtx;
   delete HostEntriesBeginToTransTable;
   delete TrlTblMtx;
   delete HostPtrToTableMap;
   delete TblMapMtx;
 }
 
 void RTLsTy::LoadRTLs() {
+
+  if (char *envStr = getenv("LIBOMPTARGET_INFO")) {
+    InfoLevel = std::stoi(envStr);
+  }
 #ifdef OMPTARGET_DEBUG
   if (char *envStr = getenv("LIBOMPTARGET_DEBUG")) {
     DebugLevel = std::stoi(envStr);
   }
 #endif // OMPTARGET_DEBUG
 
   // Parse environment variable OMP_TARGET_OFFLOAD (if set)
   TargetOffloadPolicy = (kmp_target_offload_kind_t) __kmpc_get_target_offload();

openmp/libomptarget/test/mapping/alloc_fail.c

 // RUN: %libomptarget-compile-x86_64-pc-linux-gnu
 // RUN: %libomptarget-run-fail-x86_64-pc-linux-gnu 2>&1 \
 // RUN: | %fcheck-x86_64-pc-linux-gnu
 
 // RUN: %libomptarget-compile-nvptx64-nvidia-cuda
 // RUN: %libomptarget-run-fail-nvptx64-nvidia-cuda 2>&1 \
 // RUN: | %fcheck-nvptx64-nvidia-cuda
 
+// CHECK: Libomptarget message: explicit extension not allowed: host address specified is 0x{{.*}} (8 bytes), but device allocation maps to host at 0x{{.*}} (8 bytes)
+// CHECK: Libomptarget error: Call to getOrAllocTgtPtr returned null pointer (device failure or illegal mapping).
 // CHECK: Libomptarget fatal error 1: failure of target construct while offloading is mandatory
 
 int main() {
   int arr[4] = {0, 1, 2, 3};
 #pragma omp target data map(alloc: arr[0:2])
 #pragma omp target data map(alloc: arr[1:2])
   ;
   return 0;
 }

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

 #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 error: Call to getOrAllocTgtPtr returned null pointer ('present' map type modifier).
+  // CHECK: Libomptarget error: Call to targetDataBegin failed, abort target.
+  // CHECK: Libomptarget error: Failed to process data before launching the kernel.
   // 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_array_extension.c

#pragma omp target map(present, tofrom: arr[SMALL])
     ;
   }
 
   // CHECK: arr is present
   fprintf(stderr, "arr is present\n");
 
   // CHECK: Libomptarget message: explicit extension not allowed: host address specified is 0x{{0*}}[[#LARGE_ADDR]] ([[#LARGE_BYTES]] bytes), but device allocation maps to host at 0x{{0*}}[[#SMALL_ADDR]] ([[#SMALL_BYTES]] bytes)
   // CHECK: Libomptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#LARGE_ADDR]] ([[#LARGE_BYTES]] bytes)
+  // CHECK: Libomptarget error: Call to getOrAllocTgtPtr returned null pointer ('present' map type modifier).
+  // CHECK: Libomptarget error: Call to targetDataBegin failed, abort target.
+  // CHECK: Libomptarget error: Failed to process data before launching the kernel.
   // CHECK: Libomptarget fatal error 1: failure of target construct while offloading is mandatory
 #pragma omp target data map(alloc: arr[SMALL])
   {
 #pragma omp target map(present, tofrom: arr[LARGE])
     ;
   }
 
   // CHECK-NOT: arr is present
   fprintf(stderr, "arr is present\n");
 
   return 0;
 }

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

#pragma omp target data map(present, tofrom: arr[SMALL])
     ;
   }
 
   // CHECK: arr is present
   fprintf(stderr, "arr is present\n");
 
   // CHECK: Libomptarget message: explicit extension not allowed: host address specified is 0x{{0*}}[[#LARGE_ADDR]] ([[#LARGE_BYTES]] bytes), but device allocation maps to host at 0x{{0*}}[[#SMALL_ADDR]] ([[#SMALL_BYTES]] bytes)
   // CHECK: Libomptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#LARGE_ADDR]] ([[#LARGE_BYTES]] bytes)
+  // CHECK: Libomptarget error: Call to getOrAllocTgtPtr returned null pointer ('present' map type modifier).
   // CHECK: Libomptarget fatal error 1: failure of target construct while offloading is mandatory
 #pragma omp target data map(alloc: arr[SMALL])
   {
 #pragma omp target data map(present, tofrom: arr[LARGE])
     ;
   }
 
   // CHECK-NOT: arr is present
   fprintf(stderr, "arr is present\n");
 
   return 0;
 }

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

 #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 error: Call to getOrAllocTgtPtr returned null pointer ('present' map type modifier).
   // CHECK: Libomptarget fatal error 1: failure of target construct while offloading is mandatory
 #pragma omp target enter data map(present, alloc: i)
 
   // CHECK-NOT: i is present
   fprintf(stderr, "i is present\n");
 
   return 0;
 }

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

#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 error: Call to getOrAllocTgtPtr returned null pointer ('present' map type modifier).
+  // CHECK: Libomptarget error: Call to targetDataBegin failed, abort target.
+  // CHECK: Libomptarget error: Failed to process data before launching the kernel.
   // 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;
 }