Index: openmp/libomptarget/include/omptarget.h =================================================================== --- openmp/libomptarget/include/omptarget.h +++ openmp/libomptarget/include/omptarget.h @@ -49,6 +49,8 @@ OMP_TGT_MAPTYPE_IMPLICIT = 0x200, // copy data to device OMP_TGT_MAPTYPE_CLOSE = 0x400, + // runtime error if not already allocated + OMP_TGT_MAPTYPE_PRESENT = 0x800, // member of struct, member given by [16 MSBs] - 1 OMP_TGT_MAPTYPE_MEMBER_OF = 0xffff000000000000 }; Index: openmp/libomptarget/src/device.h =================================================================== --- openmp/libomptarget/src/device.h +++ openmp/libomptarget/src/device.h @@ -177,8 +177,10 @@ uint64_t getMapEntryRefCnt(void *HstPtrBegin); LookupResult lookupMapping(void *HstPtrBegin, int64_t Size); void *getOrAllocTgtPtr(void *HstPtrBegin, void *HstPtrBase, int64_t Size, - bool &IsNew, bool &IsHostPtr, bool IsImplicit, bool UpdateRefCount = true, - bool HasCloseModifier = false); + bool &IsNew, bool &IsHostPtr, bool IsImplicit, + bool UpdateRefCount = true, + bool HasCloseModifier = false, + bool HasPresentModifier = false); void *getTgtPtrBegin(void *HstPtrBegin, int64_t Size); void *getTgtPtrBegin(void *HstPtrBegin, int64_t Size, bool &IsLast, bool UpdateRefCount, bool &IsHostPtr); Index: openmp/libomptarget/src/device.cpp =================================================================== --- openmp/libomptarget/src/device.cpp +++ openmp/libomptarget/src/device.cpp @@ -160,8 +160,10 @@ // If NULL is returned, then either data allocation failed or the user tried // to do an illegal mapping. void *DeviceTy::getOrAllocTgtPtr(void *HstPtrBegin, void *HstPtrBase, - int64_t Size, bool &IsNew, bool &IsHostPtr, bool IsImplicit, - bool UpdateRefCount, bool HasCloseModifier) { + int64_t Size, bool &IsNew, bool &IsHostPtr, + bool IsImplicit, bool UpdateRefCount, + bool HasCloseModifier, + bool HasPresentModifier) { void *rc = NULL; IsHostPtr = false; IsNew = false; @@ -190,31 +192,37 @@ } else if ((lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter) && !IsImplicit) { // Explicit extension of mapped data - not allowed. DP("Explicit extension of mapping is not allowed.\n"); - } else if (Size) { - // If unified shared memory is active, implicitly mapped variables that are not - // privatized use host address. Any explicitly mapped variables also use - // host address where correctness is not impeded. In all other cases - // maps are respected. - // In addition to the mapping rules above, the close map - // modifier forces the mapping of the variable to the device. - if (RTLs->RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY && - !HasCloseModifier) { + } else if (RTLs->RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY && + !HasCloseModifier) { + // If unified shared memory is active, implicitly mapped variables that are + // not privatized use host address. Any explicitly mapped variables also use + // host address where correctness is not impeded. In all other cases maps + // are respected. + // In addition to the mapping rules above, the close map modifier forces the + // mapping of the variable to the device. + if (Size) { DP("Return HstPtrBegin " DPxMOD " Size=%ld RefCount=%s\n", - DPxPTR((uintptr_t)HstPtrBegin), Size, (UpdateRefCount ? " updated" : "")); + DPxPTR((uintptr_t)HstPtrBegin), Size, + (UpdateRefCount ? " updated" : "")); IsHostPtr = true; rc = HstPtrBegin; - } else { - // If it is not contained and Size > 0 we should create a new entry for it. - IsNew = true; - uintptr_t tp = (uintptr_t)RTL->data_alloc(RTLDeviceID, Size, HstPtrBegin); - DP("Creating new map entry: HstBase=" DPxMOD ", HstBegin=" DPxMOD ", " - "HstEnd=" DPxMOD ", TgtBegin=" DPxMOD "\n", DPxPTR(HstPtrBase), - DPxPTR(HstPtrBegin), DPxPTR((uintptr_t)HstPtrBegin + Size), DPxPTR(tp)); - HostDataToTargetMap.emplace( - HostDataToTargetTy((uintptr_t)HstPtrBase, (uintptr_t)HstPtrBegin, - (uintptr_t)HstPtrBegin + Size, tp)); - rc = (void *)tp; } + } else if (HasPresentModifier) { + DP("Mapping required but does not exist%s for HstPtrBegin=" DPxMOD + ", Size=%ld\n", + (IsImplicit ? " (implicit)" : ""), DPxPTR(HstPtrBegin), Size); + } else if (Size) { + // If it is not contained and Size > 0, we should create a new entry for it. + IsNew = true; + uintptr_t tp = (uintptr_t)RTL->data_alloc(RTLDeviceID, Size, HstPtrBegin); + DP("Creating new map entry: HstBase=" DPxMOD ", HstBegin=" DPxMOD ", " + "HstEnd=" DPxMOD ", TgtBegin=" DPxMOD "\n", + DPxPTR(HstPtrBase), DPxPTR(HstPtrBegin), + DPxPTR((uintptr_t)HstPtrBegin + Size), DPxPTR(tp)); + HostDataToTargetMap.emplace( + HostDataToTargetTy((uintptr_t)HstPtrBase, (uintptr_t)HstPtrBegin, + (uintptr_t)HstPtrBegin + Size, tp)); + rc = (void *)tp; } DataMapMtx.unlock(); Index: openmp/libomptarget/src/omptarget.cpp =================================================================== --- openmp/libomptarget/src/omptarget.cpp +++ openmp/libomptarget/src/omptarget.cpp @@ -253,6 +253,7 @@ // Force the creation of a device side copy of the data when: // a close map modifier was associated with a map that contained a to. bool HasCloseModifier = arg_types[i] & OMP_TGT_MAPTYPE_CLOSE; + bool HasPresentModifier = arg_types[i] & OMP_TGT_MAPTYPE_PRESENT; // UpdateRef is based on MEMBER_OF instead of TARGET_PARAM because if we // have reached this point via __tgt_target_data_begin and not __tgt_target // then no argument is marked as TARGET_PARAM ("omp target data map" is not @@ -262,12 +263,25 @@ if (arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ) { DP("Has a pointer entry: \n"); // base is address of pointer. - Pointer_TgtPtrBegin = Device.getOrAllocTgtPtr(HstPtrBase, HstPtrBase, - sizeof(void *), Pointer_IsNew, IsHostPtr, IsImplicit, UpdateRef, - HasCloseModifier); + // + // The only case we can have a PTR_AND_OBJ entry is if we have a pointer + // which is member of a struct or element of an array. For example: + // + // #pragma omp target map(s.p[0:N]) + // + // The map entry for s comes first, and the PTR_AND_OBJ entry comes + // afterward, so the space for the pointer is allocated before the + // PTR_AND_OBJ entry is handled by the runtime below. As a result, + // !Pointer_TgtPtrBegin below is actually impossible and the + // HasPresentModifier argument is redundant. We keep them anyway as a + // form of defensive programming. + Pointer_TgtPtrBegin = Device.getOrAllocTgtPtr( + HstPtrBase, HstPtrBase, sizeof(void *), Pointer_IsNew, IsHostPtr, + IsImplicit, UpdateRef, HasCloseModifier, HasPresentModifier); if (!Pointer_TgtPtrBegin) { - DP("Call to getOrAllocTgtPtr returned null pointer (device failure or " - "illegal mapping).\n"); + DP("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" @@ -279,13 +293,15 @@ UpdateRef = true; // subsequently update ref count of pointee } - void *TgtPtrBegin = Device.getOrAllocTgtPtr(HstPtrBegin, HstPtrBase, - data_size, IsNew, IsHostPtr, IsImplicit, UpdateRef, HasCloseModifier); - if (!TgtPtrBegin && data_size) { - // If data_size==0, then the argument could be a zero-length pointer to - // NULL, so getOrAlloc() returning NULL is not an error. - DP("Call to getOrAllocTgtPtr returned null pointer (device failure or " - "illegal mapping).\n"); + void *TgtPtrBegin = Device.getOrAllocTgtPtr( + HstPtrBegin, HstPtrBase, data_size, IsNew, IsHostPtr, IsImplicit, + UpdateRef, HasCloseModifier, HasPresentModifier); + // If data_size==0, then the argument could be a zero-length pointer to + // NULL, so getOrAlloc() returning NULL is not an error. + if (!TgtPtrBegin && (data_size || HasPresentModifier)) { + DP("Call to getOrAllocTgtPtr returned null pointer (%s).\n", + HasPresentModifier ? "'present' map type modifier" + : "device failure or illegal mapping"); return OFFLOAD_FAIL; } DP("There are %" PRId64 " bytes allocated at target address " DPxMOD @@ -385,13 +401,21 @@ (arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ); bool ForceDelete = arg_types[i] & OMP_TGT_MAPTYPE_DELETE; bool HasCloseModifier = arg_types[i] & OMP_TGT_MAPTYPE_CLOSE; + bool HasPresentModifier = arg_types[i] & OMP_TGT_MAPTYPE_PRESENT; // If PTR_AND_OBJ, HstPtrBegin is address of pointee void *TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBegin, data_size, IsLast, UpdateRef, IsHostPtr); - DP("There are %" PRId64 " bytes allocated at target address " DPxMOD - " - is%s last\n", data_size, DPxPTR(TgtPtrBegin), - (IsLast ? "" : " not")); + if (!TgtPtrBegin && (data_size || HasPresentModifier)) { + DP("Mapping does not exist (%s)\n", + (HasPresentModifier ? "'present' map type modifier" : "ignored")); + if (HasPresentModifier) + return OFFLOAD_FAIL; + } else { + DP("There are %" PRId64 " bytes allocated at target address " DPxMOD + " - is%s last\n", + data_size, DPxPTR(TgtPtrBegin), (IsLast ? "" : " not")); + } bool DelEntry = IsLast || ForceDelete; Index: openmp/libomptarget/test/mapping/present/target.c =================================================================== --- /dev/null +++ openmp/libomptarget/test/mapping/present/target.c @@ -0,0 +1,44 @@ +// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu +// RUN: env LIBOMPTARGET_DEBUG=1 \ +// RUN: %libomptarget-run-fail-aarch64-unknown-linux-gnu 2>&1 \ +// RUN: | %fcheck-aarch64-unknown-linux-gnu + +// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu +// RUN: env LIBOMPTARGET_DEBUG=1 \ +// RUN: %libomptarget-run-fail-powerpc64-ibm-linux-gnu 2>&1 \ +// RUN: | %fcheck-powerpc64-ibm-linux-gnu + +// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu +// RUN: env LIBOMPTARGET_DEBUG=1 \ +// RUN: %libomptarget-run-fail-powerpc64le-ibm-linux-gnu 2>&1 \ +// RUN: | %fcheck-powerpc64le-ibm-linux-gnu + +// RUN: %libomptarget-compile-x86_64-pc-linux-gnu +// RUN: env LIBOMPTARGET_DEBUG=1 \ +// RUN: %libomptarget-run-fail-x86_64-pc-linux-gnu 2>&1 \ +// RUN: | %fcheck-x86_64-pc-linux-gnu + +#include + +int main() { + int i; + + // CHECK-NOT: 'present' map type modifier + // CHECK-NOT: Libomptarget fatal error +#pragma omp target data map(alloc: i) +#pragma omp target map(present, alloc: i) + ; + + // CHECK: i is present + fprintf(stderr, "i is present\n"); + + // CHECK: Libomptarget --> 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 map(present, alloc: i) + ; + + // CHECK-NOT: i is present + fprintf(stderr, "i is present\n"); + + return 0; +} Index: openmp/libomptarget/test/mapping/present/target_data.c =================================================================== --- /dev/null +++ openmp/libomptarget/test/mapping/present/target_data.c @@ -0,0 +1,44 @@ +// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu +// RUN: env LIBOMPTARGET_DEBUG=1 \ +// RUN: %libomptarget-run-fail-aarch64-unknown-linux-gnu 2>&1 \ +// RUN: | %fcheck-aarch64-unknown-linux-gnu + +// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu +// RUN: env LIBOMPTARGET_DEBUG=1 \ +// RUN: %libomptarget-run-fail-powerpc64-ibm-linux-gnu 2>&1 \ +// RUN: | %fcheck-powerpc64-ibm-linux-gnu + +// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu +// RUN: env LIBOMPTARGET_DEBUG=1 \ +// RUN: %libomptarget-run-fail-powerpc64le-ibm-linux-gnu 2>&1 \ +// RUN: | %fcheck-powerpc64le-ibm-linux-gnu + +// RUN: %libomptarget-compile-x86_64-pc-linux-gnu +// RUN: env LIBOMPTARGET_DEBUG=1 \ +// RUN: %libomptarget-run-fail-x86_64-pc-linux-gnu 2>&1 \ +// RUN: | %fcheck-x86_64-pc-linux-gnu + +#include + +int main() { + int i; + + // CHECK-NOT: 'present' map type modifier + // CHECK-NOT: Libomptarget fatal error +#pragma omp target data map(alloc: i) +#pragma omp target data map(present, alloc: i) + ; + + // CHECK: i is present + fprintf(stderr, "i is present\n"); + + // CHECK: Libomptarget --> 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(present, alloc: i) + ; + + // CHECK-NOT: i is present + fprintf(stderr, "i is present\n"); + + return 0; +} Index: openmp/libomptarget/test/mapping/present/target_enter_data.c =================================================================== --- /dev/null +++ openmp/libomptarget/test/mapping/present/target_enter_data.c @@ -0,0 +1,43 @@ +// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu +// RUN: env LIBOMPTARGET_DEBUG=1 \ +// RUN: %libomptarget-run-fail-aarch64-unknown-linux-gnu 2>&1 \ +// RUN: | %fcheck-aarch64-unknown-linux-gnu + +// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu +// RUN: env LIBOMPTARGET_DEBUG=1 \ +// RUN: %libomptarget-run-fail-powerpc64-ibm-linux-gnu 2>&1 \ +// RUN: | %fcheck-powerpc64-ibm-linux-gnu + +// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu +// RUN: env LIBOMPTARGET_DEBUG=1 \ +// RUN: %libomptarget-run-fail-powerpc64le-ibm-linux-gnu 2>&1 \ +// RUN: | %fcheck-powerpc64le-ibm-linux-gnu + +// RUN: %libomptarget-compile-x86_64-pc-linux-gnu +// RUN: env LIBOMPTARGET_DEBUG=1 \ +// RUN: %libomptarget-run-fail-x86_64-pc-linux-gnu 2>&1 \ +// RUN: | %fcheck-x86_64-pc-linux-gnu + +#include + +int main() { + int i; + + // CHECK-NOT: 'present' map type modifier + // CHECK-NOT: Libomptarget fatal error +#pragma omp target enter data map(alloc: i) +#pragma omp target enter data map(present, alloc: i) +#pragma omp target exit data map(delete: i) + + // CHECK: i is present + fprintf(stderr, "i is present\n"); + + // CHECK: Libomptarget --> Call to getOrAllocTgtPtr returned null pointer ('present' map type modifier). + // CHECK: Libomptarget fatal error 1: failure of target construct while offloading is mandatory +#pragma omp target enter data map(present, alloc: i) + + // CHECK-NOT: i is present + fprintf(stderr, "i is present\n"); + + return 0; +} Index: openmp/libomptarget/test/mapping/present/target_exit_data.c =================================================================== --- /dev/null +++ openmp/libomptarget/test/mapping/present/target_exit_data.c @@ -0,0 +1,42 @@ +// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu +// RUN: env LIBOMPTARGET_DEBUG=1 \ +// RUN: %libomptarget-run-fail-aarch64-unknown-linux-gnu 2>&1 \ +// RUN: | %fcheck-aarch64-unknown-linux-gnu + +// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu +// RUN: env LIBOMPTARGET_DEBUG=1 \ +// RUN: %libomptarget-run-fail-powerpc64-ibm-linux-gnu 2>&1 \ +// RUN: | %fcheck-powerpc64-ibm-linux-gnu + +// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu +// RUN: env LIBOMPTARGET_DEBUG=1 \ +// RUN: %libomptarget-run-fail-powerpc64le-ibm-linux-gnu 2>&1 \ +// RUN: | %fcheck-powerpc64le-ibm-linux-gnu + +// RUN: %libomptarget-compile-x86_64-pc-linux-gnu +// RUN: env LIBOMPTARGET_DEBUG=1 \ +// RUN: %libomptarget-run-fail-x86_64-pc-linux-gnu 2>&1 \ +// RUN: | %fcheck-x86_64-pc-linux-gnu + +#include + +int main() { + int i; + + // CHECK-NOT: 'present' map type modifier + // CHECK-NOT: Libomptarget fatal error +#pragma omp target enter data map(alloc: i) +#pragma omp target exit data map(present, release: i) + + // CHECK: i is present + fprintf(stderr, "i is present\n"); + + // CHECK: Libomptarget --> Mapping does not exist ('present' map type modifier) + // CHECK: Libomptarget fatal error 1: failure of target construct while offloading is mandatory +#pragma omp target exit data map(present, release: i) + + // CHECK-NOT: i is present + fprintf(stderr, "i is present\n"); + + return 0; +} Index: openmp/libomptarget/test/mapping/present/unified_shared_memory.c =================================================================== --- /dev/null +++ openmp/libomptarget/test/mapping/present/unified_shared_memory.c @@ -0,0 +1,47 @@ +// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu +// RUN: env LIBOMPTARGET_DEBUG=1 \ +// RUN: %libomptarget-run-aarch64-unknown-linux-gnu 2>&1 \ +// RUN: | %fcheck-aarch64-unknown-linux-gnu + +// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu +// RUN: env LIBOMPTARGET_DEBUG=1 \ +// RUN: %libomptarget-run-powerpc64-ibm-linux-gnu 2>&1 \ +// RUN: | %fcheck-powerpc64-ibm-linux-gnu + +// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu +// RUN: env LIBOMPTARGET_DEBUG=1 \ +// RUN: %libomptarget-run-powerpc64le-ibm-linux-gnu 2>&1 \ +// RUN: | %fcheck-powerpc64le-ibm-linux-gnu + +// RUN: %libomptarget-compile-x86_64-pc-linux-gnu +// RUN: env LIBOMPTARGET_DEBUG=1 \ +// RUN: %libomptarget-run-x86_64-pc-linux-gnu 2>&1 \ +// RUN: | %fcheck-x86_64-pc-linux-gnu + +#include + +// The runtime considers unified shared memory to be always present. +#pragma omp requires unified_shared_memory + +int main() { + int i; + + // CHECK-NOT: 'present' map type modifier + // CHECK-NOT: Libomptarget fatal error +#pragma omp target data map(alloc: i) +#pragma omp target map(present, alloc: i) + ; + + // CHECK: i is present + fprintf(stderr, "i is present\n"); + + // CHECK-NOT: 'present' map type modifier + // CHECK-NOT: Libomptarget fatal error +#pragma omp target map(present, alloc: i) + ; + + // CHECK: is present + fprintf(stderr, "i is present\n"); + + return 0; +} Index: openmp/libomptarget/test/mapping/present/zero_length_array_section.c =================================================================== --- /dev/null +++ openmp/libomptarget/test/mapping/present/zero_length_array_section.c @@ -0,0 +1,47 @@ +// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu +// RUN: env LIBOMPTARGET_DEBUG=1 \ +// RUN: %libomptarget-run-fail-aarch64-unknown-linux-gnu 2>&1 \ +// RUN: | %fcheck-aarch64-unknown-linux-gnu + +// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu +// RUN: env LIBOMPTARGET_DEBUG=1 \ +// RUN: %libomptarget-run-fail-powerpc64-ibm-linux-gnu 2>&1 \ +// RUN: | %fcheck-powerpc64-ibm-linux-gnu + +// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu +// RUN: env LIBOMPTARGET_DEBUG=1 \ +// RUN: %libomptarget-run-fail-powerpc64le-ibm-linux-gnu 2>&1 \ +// RUN: | %fcheck-powerpc64le-ibm-linux-gnu + +// RUN: %libomptarget-compile-x86_64-pc-linux-gnu +// RUN: env LIBOMPTARGET_DEBUG=1 \ +// RUN: %libomptarget-run-fail-x86_64-pc-linux-gnu 2>&1 \ +// RUN: | %fcheck-x86_64-pc-linux-gnu + +#include + +int main() { + int arr[5]; + + // CHECK-NOT: 'present' map type modifier + // CHECK-NOT: Libomptarget fatal error +#pragma omp target data map(alloc: arr[0:5]) +#pragma omp target map(present, alloc: arr[0:0]) + ; + + // CHECK: arr is present + fprintf(stderr, "arr is present\n"); + + // arr[0:0] doesn't create an actual mapping in the first directive. + // + // CHECK: Libomptarget --> 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[0:0]) +#pragma omp target map(present, alloc: arr[0:0]) + ; + + // CHECK-NOT: arr is present + fprintf(stderr, "arr is present\n"); + + return 0; +} Index: openmp/libomptarget/test/mapping/present/zero_length_array_section_exit.c =================================================================== --- /dev/null +++ openmp/libomptarget/test/mapping/present/zero_length_array_section_exit.c @@ -0,0 +1,45 @@ +// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu +// RUN: env LIBOMPTARGET_DEBUG=1 \ +// RUN: %libomptarget-run-fail-aarch64-unknown-linux-gnu 2>&1 \ +// RUN: | %fcheck-aarch64-unknown-linux-gnu + +// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu +// RUN: env LIBOMPTARGET_DEBUG=1 \ +// RUN: %libomptarget-run-fail-powerpc64-ibm-linux-gnu 2>&1 \ +// RUN: | %fcheck-powerpc64-ibm-linux-gnu + +// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu +// RUN: env LIBOMPTARGET_DEBUG=1 \ +// RUN: %libomptarget-run-fail-powerpc64le-ibm-linux-gnu 2>&1 \ +// RUN: | %fcheck-powerpc64le-ibm-linux-gnu + +// RUN: %libomptarget-compile-x86_64-pc-linux-gnu +// RUN: env LIBOMPTARGET_DEBUG=1 \ +// RUN: %libomptarget-run-fail-x86_64-pc-linux-gnu 2>&1 \ +// RUN: | %fcheck-x86_64-pc-linux-gnu + +#include + +int main() { + int arr[5]; + + // CHECK-NOT: 'present' map type modifier + // CHECK-NOT: Libomptarget fatal error +#pragma omp target enter data map(alloc: arr[0:5]) +#pragma omp target exit data map(present, release: arr[0:0]) + + // CHECK: arr is present + fprintf(stderr, "arr is present\n"); + + // arr[0:0] doesn't create an actual mapping in the first directive. + // + // CHECK: Libomptarget --> Mapping does not exist ('present' map type modifier) + // CHECK: Libomptarget fatal error 1: failure of target construct while offloading is mandatory +#pragma omp target enter data map(alloc: arr[0:0]) +#pragma omp target exit data map(present, release: arr[0:0]) + + // CHECK-NOT: arr is present + fprintf(stderr, "arr is present\n"); + + return 0; +}