[OpenMP] Always emit debug messages that indicate offloading failure

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.

Reviewers: jdoerfort

Subscribers: guansong sstefan1 yaxunl ye-luo

Tags: #OpenMP

Differential Revision: https://reviews.llvm.org/D86483

diff --git a/openmp/libomptarget/src/api.cpp b/openmp/libomptarget/src/api.cpp
index 2c5d7b5..d9258b4 100644 (file)
--- a/openmp/libomptarget/src/api.cpp
+++ b/openmp/libomptarget/src/api.cpp
@@ -132,18 +132,18 @@ EXTERN int omp_target_memcpy(void *dst, void *src, size_t length,
       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");
-      return OFFLOAD_FAIL;
+    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");
-      return OFFLOAD_FAIL;
+    REPORT("omp_target_memcpy returns OFFLOAD_FAIL\n");
+    return OFFLOAD_FAIL;
   }
 
   int rc = OFFLOAD_SUCCESS;
@@ -207,7 +207,7 @@ EXTERN int omp_target_memcpy_rect(void *dst, void *src, size_t element_size,
 
   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;
   }
 
@@ -250,17 +250,17 @@ EXTERN int omp_target_associate_ptr(void *host_ptr, void *device_ptr,
       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;
   }
 
@@ -276,17 +276,18 @@ EXTERN int omp_target_disassociate_ptr(void *host_ptr, int device_num) {
       "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;
   }
 

diff --git a/openmp/libomptarget/src/device.cpp b/openmp/libomptarget/src/device.cpp
index cca0123..fdf625c 100644 (file)
--- a/openmp/libomptarget/src/device.cpp
+++ b/openmp/libomptarget/src/device.cpp
@@ -67,8 +67,8 @@ int DeviceTy::associatePtr(void *HstPtrBegin, void *TgtPtrBegin, int64_t Size) {
          "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;
     }
   }
@@ -103,14 +103,14 @@ int DeviceTy::disassociatePtr(void *HstPtrBegin) {
       DataMapMtx.unlock();
       return OFFLOAD_SUCCESS;
     } else {
-      DP("Trying to disassociate a pointer which was not mapped via "
-         "omp_target_associate_ptr\n");
+      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;
 }
 
@@ -348,8 +348,9 @@ int DeviceTy::deallocTgtPtr(void *HstPtrBegin, int64_t Size, bool ForceDelete,
     }
     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;
   }
 

diff --git a/openmp/libomptarget/src/interface.cpp b/openmp/libomptarget/src/interface.cpp
index 24b733c..d15e4f3 100644 (file)
--- a/openmp/libomptarget/src/interface.cpp
+++ b/openmp/libomptarget/src/interface.cpp
@@ -62,6 +62,8 @@ static void HandleTargetOutcome(bool success) {
       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;
@@ -303,7 +305,7 @@ EXTERN int __tgt_target_mapper(int64_t device_id, void *host_ptr,
   }
 
   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;
   }
@@ -363,7 +365,7 @@ EXTERN int __tgt_target_teams_mapper(int64_t device_id, void *host_ptr,
   }
 
   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;
   }

diff --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp
index 9b4b85d..e20a2c9 100644 (file)
--- a/openmp/libomptarget/src/omptarget.cpp
+++ b/openmp/libomptarget/src/omptarget.cpp
@@ -23,8 +23,7 @@
 #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.
@@ -87,7 +86,7 @@ static int InitLibrary(DeviceTy& Device) {
            "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;
     }
@@ -96,7 +95,7 @@ static int InitLibrary(DeviceTy& Device) {
         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;
@@ -109,8 +108,8 @@ static int InitLibrary(DeviceTy& Device) {
 
     // Invalid image for these host entries!
     if (hsize != tsize) {
-      DP("Host and Target tables mismatch for device id %d [%zx != %zx].\n",
-         device_id, hsize, tsize);
+      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;
@@ -169,7 +168,7 @@ static int InitLibrary(DeviceTy& Device) {
           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;
           }
@@ -191,7 +190,7 @@ static int InitLibrary(DeviceTy& Device) {
 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;
   }
 
@@ -203,7 +202,7 @@ int CheckDeviceAndCtors(int64_t device_id) {
   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;
   }
 
@@ -275,8 +274,8 @@ int targetDataBegin(DeviceTy &Device, int32_t arg_num, void **args_base,
                                 arg_types[i], arg_mappers[i], targetDataBegin);
 
       if (rc != OFFLOAD_SUCCESS) {
-        DP("Call to targetDataBegin via targetDataMapper for custom mapper"
-           " failed.\n");
+        REPORT("Call to targetDataBegin via targetDataMapper for custom mapper"
+               " failed.\n");
         return OFFLOAD_FAIL;
       }
 
@@ -338,9 +337,9 @@ int targetDataBegin(DeviceTy &Device, int32_t arg_num, void **args_base,
           HstPtrBase, HstPtrBase, sizeof(void *), Pointer_IsNew, IsHostPtr,
           IsImplicit, UpdateRef, HasCloseModifier, HasPresentModifier);
       if (!PointerTgtPtrBegin) {
-        DP("Call to getOrAllocTgtPtr returned null pointer (%s).\n",
-           HasPresentModifier ? "'present' map type modifier"
-                              : "device failure or illegal mapping");
+        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"
@@ -358,9 +357,9 @@ int targetDataBegin(DeviceTy &Device, int32_t arg_num, void **args_base,
     // 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",
-         HasPresentModifier ? "'present' map type modifier"
-                            : "device failure or illegal mapping");
+      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
@@ -397,7 +396,7 @@ int targetDataBegin(DeviceTy &Device, int32_t arg_num, void **args_base,
         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;
         }
       }
@@ -411,7 +410,7 @@ int targetDataBegin(DeviceTy &Device, int32_t arg_num, void **args_base,
       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
@@ -469,8 +468,8 @@ int targetDataEnd(DeviceTy &Device, int32_t ArgNum, void **ArgBases,
                              ArgTypes[I], ArgMappers[I], targetDataEnd);
 
       if (Ret != OFFLOAD_SUCCESS) {
-        DP("Call to targetDataEnd via targetDataMapper for custom mapper"
-           " failed.\n");
+        REPORT("Call to targetDataEnd via targetDataMapper for custom mapper"
+               " failed.\n");
         return OFFLOAD_FAIL;
       }
 
@@ -563,7 +562,7 @@ int targetDataEnd(DeviceTy &Device, int32_t ArgNum, void **ArgBases,
           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;
           }
         }
@@ -622,7 +621,7 @@ int targetDataEnd(DeviceTy &Device, int32_t ArgNum, void **ArgBases,
   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;
     }
   }
@@ -632,7 +631,7 @@ int targetDataEnd(DeviceTy &Device, int32_t ArgNum, void **ArgBases,
     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;
     }
   }
@@ -663,8 +662,9 @@ int target_data_update(DeviceTy &Device, int32_t arg_num,
                            arg_types[i], arg_mappers[i], target_data_update);
 
       if (rc != OFFLOAD_SUCCESS) {
-        DP("Call to target_data_update via targetDataMapper for custom mapper"
-           " failed.\n");
+        REPORT(
+            "Call to target_data_update via targetDataMapper for custom mapper"
+            " failed.\n");
         return OFFLOAD_FAIL;
       }
 
@@ -700,7 +700,7 @@ int target_data_update(DeviceTy &Device, int32_t arg_num,
           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;
       }
 
@@ -727,7 +727,7 @@ int target_data_update(DeviceTy &Device, int32_t arg_num,
           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;
       }
 
@@ -747,7 +747,7 @@ int target_data_update(DeviceTy &Device, int32_t arg_num,
         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;
         }
@@ -997,7 +997,7 @@ int processDataBefore(int64_t DeviceId, void *HostPtr, int32_t ArgNum,
   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;
   }
 
@@ -1044,7 +1044,7 @@ int processDataBefore(int64_t DeviceId, void *HostPtr, int32_t ArgNum,
         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;
         }
       }
@@ -1067,8 +1067,8 @@ int processDataBefore(int64_t DeviceId, void *HostPtr, int32_t ArgNum,
                                           TgtBaseOffset, IsFirstPrivate,
                                           TgtPtrBegin, TgtArgs.size());
       if (Ret != OFFLOAD_SUCCESS) {
-        DP("Failed to process %sprivate argument " DPxMOD "\n",
-           (IsFirstPrivate ? "first-" : ""), DPxPTR(HstPtrBegin));
+        REPORT("Failed to process %sprivate argument " DPxMOD "\n",
+               (IsFirstPrivate ? "first-" : ""), DPxPTR(HstPtrBegin));
         return OFFLOAD_FAIL;
       }
     } else {
@@ -1114,14 +1114,14 @@ int processDataAfter(int64_t DeviceId, void *HostPtr, int32_t ArgNum,
   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;
   }
 
@@ -1143,8 +1143,8 @@ int target(int64_t DeviceId, void *HostPtr, int32_t ArgNum, void **ArgBases,
   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",
-       DPxPTR(HostPtr));
+    REPORT("Host ptr " DPxMOD " does not have a matching target pointer.\n",
+           DPxPTR(HostPtr));
     return OFFLOAD_FAIL;
   }
 
@@ -1170,7 +1170,7 @@ int target(int64_t DeviceId, void *HostPtr, int32_t ArgNum, void **ArgBases,
                               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;
   }
 
@@ -1191,7 +1191,7 @@ int target(int64_t DeviceId, void *HostPtr, int32_t ArgNum, void **ArgBases,
                            TgtArgs.size(), &AsyncInfo);
 
   if (Ret != OFFLOAD_SUCCESS) {
-    DP("Executing target region abort target.\n");
+    REPORT("Executing target region abort target.\n");
     return OFFLOAD_FAIL;
   }
 
@@ -1201,7 +1201,7 @@ int target(int64_t DeviceId, void *HostPtr, int32_t ArgNum, void **ArgBases,
                          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;
   }
 

diff --git a/openmp/libomptarget/src/private.h b/openmp/libomptarget/src/private.h
index c5ca21d..a7091c5 100644 (file)
--- a/openmp/libomptarget/src/private.h
+++ b/openmp/libomptarget/src/private.h
@@ -106,6 +106,12 @@ typedef int (*TargetDataFuncPtrTy)(DeviceTy &, int32_t, void **, void **,
     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" {
@@ -119,6 +125,7 @@ int __kmpc_get_target_offload(void) __attribute__((weak));
 }
 #endif
 
+extern int InfoLevel;
 #ifdef OMPTARGET_DEBUG
 extern int DebugLevel;
 
@@ -132,4 +139,18 @@ extern int DebugLevel;
 #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

diff --git a/openmp/libomptarget/src/rtl.cpp b/openmp/libomptarget/src/rtl.cpp
index ed339e0..620451b 100644 (file)
--- a/openmp/libomptarget/src/rtl.cpp
+++ b/openmp/libomptarget/src/rtl.cpp
@@ -61,6 +61,10 @@ __attribute__((destructor(101))) void deinit() {
 }
 
 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);

diff --git a/openmp/libomptarget/test/mapping/alloc_fail.c b/openmp/libomptarget/test/mapping/alloc_fail.c
index 6d1f708..2566796 100644 (file)
--- a/openmp/libomptarget/test/mapping/alloc_fail.c
+++ b/openmp/libomptarget/test/mapping/alloc_fail.c
@@ -18,6 +18,8 @@
 // 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() {

diff --git a/openmp/libomptarget/test/mapping/present/target.c b/openmp/libomptarget/test/mapping/present/target.c
index 1d61dc0..b37eca2 100644 (file)
--- a/openmp/libomptarget/test/mapping/present/target.c
+++ b/openmp/libomptarget/test/mapping/present/target.c
@@ -31,6 +31,9 @@ int main() {
   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)
   ;

diff --git a/openmp/libomptarget/test/mapping/present/target_array_extension.c b/openmp/libomptarget/test/mapping/present/target_array_extension.c
index 870be39..e02319c 100644 (file)
--- a/openmp/libomptarget/test/mapping/present/target_array_extension.c
+++ b/openmp/libomptarget/test/mapping/present/target_array_extension.c
@@ -98,6 +98,9 @@ int main() {
 
   // 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])
   {

diff --git a/openmp/libomptarget/test/mapping/present/target_data_array_extension.c b/openmp/libomptarget/test/mapping/present/target_data_array_extension.c
index 3aef777..b5e43c2 100644 (file)
--- a/openmp/libomptarget/test/mapping/present/target_data_array_extension.c
+++ b/openmp/libomptarget/test/mapping/present/target_data_array_extension.c
@@ -98,6 +98,7 @@ int main() {
 
   // 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])
   {

diff --git a/openmp/libomptarget/test/mapping/present/target_enter_data.c b/openmp/libomptarget/test/mapping/present/target_enter_data.c
index d96e7a4..dfe54ff 100644 (file)
--- a/openmp/libomptarget/test/mapping/present/target_enter_data.c
+++ b/openmp/libomptarget/test/mapping/present/target_enter_data.c
@@ -31,6 +31,7 @@ int main() {
   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)
 

diff --git a/openmp/libomptarget/test/mapping/present/zero_length_array_section.c b/openmp/libomptarget/test/mapping/present/zero_length_array_section.c
index 5488888..3da0ab4 100644 (file)
--- a/openmp/libomptarget/test/mapping/present/zero_length_array_section.c
+++ b/openmp/libomptarget/test/mapping/present/zero_length_array_section.c
@@ -33,6 +33,9 @@ int main() {
   // 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])