From fc247c8f3c61c327f58ba361ab01ce72641bbdcb Mon Sep 17 00:00:00 2001 From: "Joel E. Denny" Date: Wed, 22 Jul 2020 11:22:08 -0400 Subject: [PATCH] Revert "[OpenMP] Implement TR8 `present` map type modifier in runtime (2/2)" This reverts commit 45b8f7ec35ef653bafdf48034857222517c17781. It attempts to use debug macros `DPxMOD` and `DPxPTR` in release builds. Will fix and reapply later. --- openmp/libomptarget/include/omptarget.h | 2 - openmp/libomptarget/src/device.cpp | 57 ++++++++------------ openmp/libomptarget/src/device.h | 5 +- openmp/libomptarget/src/omptarget.cpp | 62 ++++++---------------- openmp/libomptarget/src/private.h | 12 +---- openmp/libomptarget/test/mapping/present/target.c | 42 --------------- .../test/mapping/present/target_data.c | 42 --------------- .../test/mapping/present/target_enter_data.c | 41 -------------- .../test/mapping/present/target_exit_data.c | 40 -------------- .../test/mapping/present/unified_shared_memory.c | 41 -------------- .../mapping/present/zero_length_array_section.c | 45 ---------------- .../present/zero_length_array_section_exit.c | 43 --------------- 12 files changed, 42 insertions(+), 390 deletions(-) delete mode 100644 openmp/libomptarget/test/mapping/present/target.c delete mode 100644 openmp/libomptarget/test/mapping/present/target_data.c delete mode 100644 openmp/libomptarget/test/mapping/present/target_enter_data.c delete mode 100644 openmp/libomptarget/test/mapping/present/target_exit_data.c delete mode 100644 openmp/libomptarget/test/mapping/present/unified_shared_memory.c delete mode 100644 openmp/libomptarget/test/mapping/present/zero_length_array_section.c delete mode 100644 openmp/libomptarget/test/mapping/present/zero_length_array_section_exit.c diff --git a/openmp/libomptarget/include/omptarget.h b/openmp/libomptarget/include/omptarget.h index e37449a..95d7158 100644 --- a/openmp/libomptarget/include/omptarget.h +++ b/openmp/libomptarget/include/omptarget.h @@ -49,8 +49,6 @@ enum tgt_map_type { OMP_TGT_MAPTYPE_IMPLICIT = 0x200, // copy data to device OMP_TGT_MAPTYPE_CLOSE = 0x400, - // runtime error if not already allocated - OMP_TGT_MAPTYPE_PRESENT = 0x1000, // member of struct, member given by [16 MSBs] - 1 OMP_TGT_MAPTYPE_MEMBER_OF = 0xffff000000000000 }; diff --git a/openmp/libomptarget/src/device.cpp b/openmp/libomptarget/src/device.cpp index 867083f..5753543 100644 --- a/openmp/libomptarget/src/device.cpp +++ b/openmp/libomptarget/src/device.cpp @@ -160,10 +160,8 @@ LookupResult DeviceTy::lookupMapping(void *HstPtrBegin, int64_t Size) { // 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, - bool HasPresentModifier) { + int64_t Size, bool &IsNew, bool &IsHostPtr, bool IsImplicit, + bool UpdateRefCount, bool HasCloseModifier) { void *rc = NULL; IsHostPtr = false; IsNew = false; @@ -192,40 +190,31 @@ void *DeviceTy::getOrAllocTgtPtr(void *HstPtrBegin, void *HstPtrBase, } 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 (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) { + } 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) { 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 by 'present' map type modifier does not exist for " - "HstPtrBegin=" DPxMOD ", Size=%ld\n", - DPxPTR(HstPtrBegin), Size); - MESSAGE("device mapping required by 'present' map type modifier does not " - "exist for host address " DPxMOD " (%ld bytes)", - 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(); diff --git a/openmp/libomptarget/src/device.h b/openmp/libomptarget/src/device.h index ebec76c..309785a 100644 --- a/openmp/libomptarget/src/device.h +++ b/openmp/libomptarget/src/device.h @@ -177,9 +177,8 @@ struct DeviceTy { 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, bool HasCloseModifier, - bool HasPresentModifier); + bool &IsNew, bool &IsHostPtr, bool IsImplicit, bool UpdateRefCount = true, + bool HasCloseModifier = false); void *getTgtPtrBegin(void *HstPtrBegin, int64_t Size); void *getTgtPtrBegin(void *HstPtrBegin, int64_t Size, bool &IsLast, bool UpdateRefCount, bool &IsHostPtr); diff --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp index 47971b9..6b4549b 100644 --- a/openmp/libomptarget/src/omptarget.cpp +++ b/openmp/libomptarget/src/omptarget.cpp @@ -308,7 +308,6 @@ int target_data_begin(DeviceTy &Device, int32_t arg_num, void **args_base, // 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 @@ -317,26 +316,13 @@ int target_data_begin(DeviceTy &Device, int32_t arg_num, void **args_base, bool UpdateRef = !(arg_types[i] & OMP_TGT_MAPTYPE_MEMBER_OF); if (arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ) { DP("Has a pointer entry: \n"); - // Base is address of pointer. - // - // Usually, the pointer is already allocated by this time. 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 pointer is already allocated by the time the - // PTR_AND_OBJ entry is handled below, and Pointer_TgtPtrBegin is thus - // non-null. However, "declare target link" can produce a PTR_AND_OBJ - // entry for a global that might not already be allocated by the time the - // PTR_AND_OBJ entry is handled below, and so the allocation might fail - // when HasPresentModifier. - Pointer_TgtPtrBegin = Device.getOrAllocTgtPtr( - HstPtrBase, HstPtrBase, sizeof(void *), Pointer_IsNew, IsHostPtr, - IsImplicit, UpdateRef, HasCloseModifier, HasPresentModifier); + // base is address of pointer. + Pointer_TgtPtrBegin = Device.getOrAllocTgtPtr(HstPtrBase, HstPtrBase, + sizeof(void *), Pointer_IsNew, IsHostPtr, IsImplicit, UpdateRef, + HasCloseModifier); if (!Pointer_TgtPtrBegin) { - DP("Call to getOrAllocTgtPtr returned null pointer (%s).\n", - HasPresentModifier ? "'present' map type modifier" - : "device failure or illegal mapping"); + DP("Call to getOrAllocTgtPtr returned null pointer (device failure or " + "illegal mapping).\n"); return OFFLOAD_FAIL; } DP("There are %zu bytes allocated at target address " DPxMOD " - is%s new" @@ -348,15 +334,13 @@ int target_data_begin(DeviceTy &Device, int32_t arg_num, void **args_base, UpdateRef = true; // subsequently update ref count of pointee } - 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"); + 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"); return OFFLOAD_FAIL; } DP("There are %" PRId64 " bytes allocated at target address " DPxMOD @@ -475,27 +459,13 @@ int target_data_end(DeviceTy &Device, int32_t arg_num, void **args_base, (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); - if (!TgtPtrBegin && (data_size || HasPresentModifier)) { - DP("Mapping does not exist (%s)\n", - (HasPresentModifier ? "'present' map type modifier" : "ignored")); - if (HasPresentModifier) { - // FIXME: This should not be an error on exit from "omp target data", - // but it should be an error upon entering an "omp target exit data". - MESSAGE("device mapping required by 'present' map type modifier does " - "not exist for host address " DPxMOD " (%ld bytes)", - DPxPTR(HstPtrBegin), data_size); - return OFFLOAD_FAIL; - } - } else { - DP("There are %" PRId64 " bytes allocated at target address " DPxMOD - " - is%s last\n", - data_size, DPxPTR(TgtPtrBegin), (IsLast ? "" : " not")); - } + DP("There are %" PRId64 " bytes allocated at target address " DPxMOD + " - is%s last\n", data_size, DPxPTR(TgtPtrBegin), + (IsLast ? "" : " not")); bool DelEntry = IsLast || ForceDelete; diff --git a/openmp/libomptarget/src/private.h b/openmp/libomptarget/src/private.h index 7772175..cb20d8c 100644 --- a/openmp/libomptarget/src/private.h +++ b/openmp/libomptarget/src/private.h @@ -80,19 +80,9 @@ typedef int (*TargetDataFuncPtrTy)(DeviceTy &, int32_t, void **, void **, int64_t *, int64_t *, void **, __tgt_async_info *); //////////////////////////////////////////////////////////////////////////////// -// implementation for messages +// implementation for fatal messages //////////////////////////////////////////////////////////////////////////////// -#define MESSAGE0(_str) \ - do { \ - fprintf(stderr, "Libomptarget message: %s\n", _str); \ - } while (0) - -#define MESSAGE(_str, ...) \ - do { \ - fprintf(stderr, "Libomptarget message: " _str "\n", __VA_ARGS__); \ - } while (0) - #define FATAL_MESSAGE0(_num, _str) \ do { \ fprintf(stderr, "Libomptarget fatal error %d: %s\n", _num, _str); \ diff --git a/openmp/libomptarget/test/mapping/present/target.c b/openmp/libomptarget/test/mapping/present/target.c deleted file mode 100644 index 1d61dc0..0000000 --- a/openmp/libomptarget/test/mapping/present/target.c +++ /dev/null @@ -1,42 +0,0 @@ -// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu -fopenmp-version=51 -// RUN: %libomptarget-run-fail-aarch64-unknown-linux-gnu 2>&1 \ -// RUN: | %fcheck-aarch64-unknown-linux-gnu - -// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu -fopenmp-version=51 -// RUN: %libomptarget-run-fail-powerpc64-ibm-linux-gnu 2>&1 \ -// RUN: | %fcheck-powerpc64-ibm-linux-gnu - -// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu -fopenmp-version=51 -// RUN: %libomptarget-run-fail-powerpc64le-ibm-linux-gnu 2>&1 \ -// RUN: | %fcheck-powerpc64le-ibm-linux-gnu - -// RUN: %libomptarget-compile-x86_64-pc-linux-gnu -fopenmp-version=51 -// RUN: %libomptarget-run-fail-x86_64-pc-linux-gnu 2>&1 \ -// RUN: | %fcheck-x86_64-pc-linux-gnu - -#include - -int main() { - int i; - - // CHECK: addr=0x[[#%x,HOST_ADDR:]], size=[[#%u,SIZE:]] - fprintf(stderr, "addr=%p, size=%ld\n", &i, sizeof i); - - // CHECK-NOT: Libomptarget -#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 message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]] bytes) - // 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; -} diff --git a/openmp/libomptarget/test/mapping/present/target_data.c b/openmp/libomptarget/test/mapping/present/target_data.c deleted file mode 100644 index fd3107d..0000000 --- a/openmp/libomptarget/test/mapping/present/target_data.c +++ /dev/null @@ -1,42 +0,0 @@ -// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu -fopenmp-version=51 -// RUN: %libomptarget-run-fail-aarch64-unknown-linux-gnu 2>&1 \ -// RUN: | %fcheck-aarch64-unknown-linux-gnu - -// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu -fopenmp-version=51 -// RUN: %libomptarget-run-fail-powerpc64-ibm-linux-gnu 2>&1 \ -// RUN: | %fcheck-powerpc64-ibm-linux-gnu - -// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu -fopenmp-version=51 -// RUN: %libomptarget-run-fail-powerpc64le-ibm-linux-gnu 2>&1 \ -// RUN: | %fcheck-powerpc64le-ibm-linux-gnu - -// RUN: %libomptarget-compile-x86_64-pc-linux-gnu -fopenmp-version=51 -// RUN: %libomptarget-run-fail-x86_64-pc-linux-gnu 2>&1 \ -// RUN: | %fcheck-x86_64-pc-linux-gnu - -#include - -int main() { - int i; - - // CHECK: addr=0x[[#%x,HOST_ADDR:]], size=[[#%u,SIZE:]] - fprintf(stderr, "addr=%p, size=%ld\n", &i, sizeof i); - - // CHECK-NOT: Libomptarget -#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 message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]] bytes) - // 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; -} diff --git a/openmp/libomptarget/test/mapping/present/target_enter_data.c b/openmp/libomptarget/test/mapping/present/target_enter_data.c deleted file mode 100644 index d96e7a4..0000000 --- a/openmp/libomptarget/test/mapping/present/target_enter_data.c +++ /dev/null @@ -1,41 +0,0 @@ -// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu -fopenmp-version=51 -// RUN: %libomptarget-run-fail-aarch64-unknown-linux-gnu 2>&1 \ -// RUN: | %fcheck-aarch64-unknown-linux-gnu - -// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu -fopenmp-version=51 -// RUN: %libomptarget-run-fail-powerpc64-ibm-linux-gnu 2>&1 \ -// RUN: | %fcheck-powerpc64-ibm-linux-gnu - -// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu -fopenmp-version=51 -// RUN: %libomptarget-run-fail-powerpc64le-ibm-linux-gnu 2>&1 \ -// RUN: | %fcheck-powerpc64le-ibm-linux-gnu - -// RUN: %libomptarget-compile-x86_64-pc-linux-gnu -fopenmp-version=51 -// RUN: %libomptarget-run-fail-x86_64-pc-linux-gnu 2>&1 \ -// RUN: | %fcheck-x86_64-pc-linux-gnu - -#include - -int main() { - int i; - - // CHECK: addr=0x[[#%x,HOST_ADDR:]], size=[[#%u,SIZE:]] - fprintf(stderr, "addr=%p, size=%ld\n", &i, sizeof i); - - // CHECK-NOT: Libomptarget -#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 message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]] bytes) - // 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; -} diff --git a/openmp/libomptarget/test/mapping/present/target_exit_data.c b/openmp/libomptarget/test/mapping/present/target_exit_data.c deleted file mode 100644 index 86b7ad8..0000000 --- a/openmp/libomptarget/test/mapping/present/target_exit_data.c +++ /dev/null @@ -1,40 +0,0 @@ -// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu -fopenmp-version=51 -// RUN: %libomptarget-run-fail-aarch64-unknown-linux-gnu 2>&1 \ -// RUN: | %fcheck-aarch64-unknown-linux-gnu - -// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu -fopenmp-version=51 -// RUN: %libomptarget-run-fail-powerpc64-ibm-linux-gnu 2>&1 \ -// RUN: | %fcheck-powerpc64-ibm-linux-gnu - -// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu -fopenmp-version=51 -// RUN: %libomptarget-run-fail-powerpc64le-ibm-linux-gnu 2>&1 \ -// RUN: | %fcheck-powerpc64le-ibm-linux-gnu - -// RUN: %libomptarget-compile-x86_64-pc-linux-gnu -fopenmp-version=51 -// RUN: %libomptarget-run-fail-x86_64-pc-linux-gnu 2>&1 \ -// RUN: | %fcheck-x86_64-pc-linux-gnu - -#include - -int main() { - int i; - - // CHECK: addr=0x[[#%x,HOST_ADDR:]], size=[[#%u,SIZE:]] - fprintf(stderr, "addr=%p, size=%ld\n", &i, sizeof i); - - // CHECK-NOT: Libomptarget -#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 message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]] bytes) - // 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; -} diff --git a/openmp/libomptarget/test/mapping/present/unified_shared_memory.c b/openmp/libomptarget/test/mapping/present/unified_shared_memory.c deleted file mode 100644 index 22d8746..0000000 --- a/openmp/libomptarget/test/mapping/present/unified_shared_memory.c +++ /dev/null @@ -1,41 +0,0 @@ -// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu -fopenmp-version=51 -// RUN: %libomptarget-run-aarch64-unknown-linux-gnu 2>&1 \ -// RUN: | %fcheck-aarch64-unknown-linux-gnu - -// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu -fopenmp-version=51 -// RUN: %libomptarget-run-powerpc64-ibm-linux-gnu 2>&1 \ -// RUN: | %fcheck-powerpc64-ibm-linux-gnu - -// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu -fopenmp-version=51 -// RUN: %libomptarget-run-powerpc64le-ibm-linux-gnu 2>&1 \ -// RUN: | %fcheck-powerpc64le-ibm-linux-gnu - -// RUN: %libomptarget-compile-x86_64-pc-linux-gnu -fopenmp-version=51 -// 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: Libomptarget -#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: Libomptarget -#pragma omp target map(present, alloc: i) - ; - - // CHECK: is present - fprintf(stderr, "i is present\n"); - - return 0; -} diff --git a/openmp/libomptarget/test/mapping/present/zero_length_array_section.c b/openmp/libomptarget/test/mapping/present/zero_length_array_section.c deleted file mode 100644 index 5488888..0000000 --- a/openmp/libomptarget/test/mapping/present/zero_length_array_section.c +++ /dev/null @@ -1,45 +0,0 @@ -// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu -fopenmp-version=51 -// RUN: %libomptarget-run-fail-aarch64-unknown-linux-gnu 2>&1 \ -// RUN: | %fcheck-aarch64-unknown-linux-gnu - -// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu -fopenmp-version=51 -// RUN: %libomptarget-run-fail-powerpc64-ibm-linux-gnu 2>&1 \ -// RUN: | %fcheck-powerpc64-ibm-linux-gnu - -// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu -fopenmp-version=51 -// RUN: %libomptarget-run-fail-powerpc64le-ibm-linux-gnu 2>&1 \ -// RUN: | %fcheck-powerpc64le-ibm-linux-gnu - -// RUN: %libomptarget-compile-x86_64-pc-linux-gnu -fopenmp-version=51 -// 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: addr=0x[[#%x,HOST_ADDR:]] - fprintf(stderr, "addr=%p\n", arr); - - // CHECK-NOT: Libomptarget -#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 message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] (0 bytes) - // 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; -} diff --git a/openmp/libomptarget/test/mapping/present/zero_length_array_section_exit.c b/openmp/libomptarget/test/mapping/present/zero_length_array_section_exit.c deleted file mode 100644 index bedc6a2..0000000 --- a/openmp/libomptarget/test/mapping/present/zero_length_array_section_exit.c +++ /dev/null @@ -1,43 +0,0 @@ -// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu -fopenmp-version=51 -// RUN: %libomptarget-run-fail-aarch64-unknown-linux-gnu 2>&1 \ -// RUN: | %fcheck-aarch64-unknown-linux-gnu - -// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu -fopenmp-version=51 -// RUN: %libomptarget-run-fail-powerpc64-ibm-linux-gnu 2>&1 \ -// RUN: | %fcheck-powerpc64-ibm-linux-gnu - -// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu -fopenmp-version=51 -// RUN: %libomptarget-run-fail-powerpc64le-ibm-linux-gnu 2>&1 \ -// RUN: | %fcheck-powerpc64le-ibm-linux-gnu - -// RUN: %libomptarget-compile-x86_64-pc-linux-gnu -fopenmp-version=51 -// 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: addr=0x[[#%x,HOST_ADDR:]] - fprintf(stderr, "addr=%p\n", arr); - - // CHECK-NOT: Libomptarget -#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 message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] (0 bytes) - // 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; -} -- 2.7.4