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])