diff --git a/openmp/libomptarget/src/api.cpp b/openmp/libomptarget/src/api.cpp
--- a/openmp/libomptarget/src/api.cpp
+++ b/openmp/libomptarget/src/api.cpp
@@ -132,18 +132,18 @@
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 @@
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 @@
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 @@
"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
--- a/openmp/libomptarget/src/device.cpp
+++ b/openmp/libomptarget/src/device.cpp
@@ -67,8 +67,8 @@
"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 @@
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 @@
}
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
--- a/openmp/libomptarget/src/interface.cpp
+++ b/openmp/libomptarget/src/interface.cpp
@@ -62,6 +62,8 @@
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 @@
}
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 @@
}
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
--- 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 @@
"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 @@
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 @@
// 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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
// 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 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 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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 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 @@
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 @@
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 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 @@
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 @@
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 @@
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 @@
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
--- a/openmp/libomptarget/src/private.h
+++ b/openmp/libomptarget/src/private.h
@@ -106,6 +106,12 @@
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 @@
}
#endif
+extern int InfoLevel;
#ifdef OMPTARGET_DEBUG
extern int DebugLevel;
@@ -132,4 +139,18 @@
#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
--- a/openmp/libomptarget/src/rtl.cpp
+++ b/openmp/libomptarget/src/rtl.cpp
@@ -61,6 +61,10 @@
}
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
--- 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
--- a/openmp/libomptarget/test/mapping/present/target.c
+++ b/openmp/libomptarget/test/mapping/present/target.c
@@ -31,6 +31,9 @@
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
--- a/openmp/libomptarget/test/mapping/present/target_array_extension.c
+++ b/openmp/libomptarget/test/mapping/present/target_array_extension.c
@@ -98,6 +98,9 @@
// 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
--- 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 @@
// 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
--- a/openmp/libomptarget/test/mapping/present/target_enter_data.c
+++ b/openmp/libomptarget/test/mapping/present/target_enter_data.c
@@ -31,6 +31,7 @@
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
--- 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 @@
// 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])