[Openmp-commits] [openmp] 7a5a74e - [OpenMP] Always emit debug messages that indicate offloading failure
Joseph Huber via Openmp-commits
openmp-commits at lists.llvm.org
Wed Aug 26 16:56:16 PDT 2020
Author: Joseph Huber
Date: 2020-08-26T19:30:41-04:00
New Revision: 7a5a74ea9675008589593e0f811c9b60fc962d0b
URL: https://github.com/llvm/llvm-project/commit/7a5a74ea9675008589593e0f811c9b60fc962d0b
DIFF: https://github.com/llvm/llvm-project/commit/7a5a74ea9675008589593e0f811c9b60fc962d0b.diff
LOG: [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
Added:
Modified:
openmp/libomptarget/src/api.cpp
openmp/libomptarget/src/device.cpp
openmp/libomptarget/src/interface.cpp
openmp/libomptarget/src/omptarget.cpp
openmp/libomptarget/src/private.h
openmp/libomptarget/src/rtl.cpp
openmp/libomptarget/test/mapping/alloc_fail.c
openmp/libomptarget/test/mapping/present/target.c
openmp/libomptarget/test/mapping/present/target_array_extension.c
openmp/libomptarget/test/mapping/present/target_data_array_extension.c
openmp/libomptarget/test/mapping/present/target_enter_data.c
openmp/libomptarget/test/mapping/present/zero_length_array_section.c
Removed:
################################################################################
diff --git a/openmp/libomptarget/src/api.cpp b/openmp/libomptarget/src/api.cpp
index 2c5d7b5ceba7..d9258b4bff5d 100644
--- 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 cca0123465ab..fdf625cb71f6 100644
--- 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
diff erent device ptr+offset with the "
- "same host ptr\n");
+ REPORT("Not allowed to re-associate a
diff erent 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 24b733ca99d2..d15e4f321fc9 100644
--- 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 9b4b85dde325..e20a2c9d4b24 100644
--- 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 c5ca21d12e5e..a7091c5b746c 100644
--- 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 ed339e034501..620451baaa0b 100644
--- 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 6d1f708dcb2e..256679606edc 100644
--- 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 1d61dc06baa4..b37eca2bc2bf 100644
--- 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 870be3956c46..e02319c98693 100644
--- 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 3aef7772f9e2..b5e43c2e1ddc 100644
--- 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 d96e7a4140da..dfe54ffbb5f2 100644
--- 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 5488888e1401..3da0ab49b61e 100644
--- 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])
More information about the Openmp-commits
mailing list