From: Joel E. Denny Date: Wed, 22 Jul 2020 18:04:58 +0000 (-0400) Subject: [OpenMP] Implement TR8 `present` map type modifier in runtime (2/2) X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=708752b2f6c55eec85accf4d67b9e9da5a08ddf1;p=platform%2Fupstream%2Fllvm.git [OpenMP] Implement TR8 `present` map type modifier in runtime (2/2) This implements OpenMP runtime support for the OpenMP TR8 `present` map type modifier. The previous patch in this series implements Clang front end support. See that patch summary for behaviors that are not yet supported. Reviewed By: grokos, jdoerfert Differential Revision: https://reviews.llvm.org/D83062 --- diff --git a/openmp/libomptarget/include/omptarget.h b/openmp/libomptarget/include/omptarget.h index 95d7158..0751816 100644 --- a/openmp/libomptarget/include/omptarget.h +++ b/openmp/libomptarget/include/omptarget.h @@ -49,6 +49,8 @@ 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 }; @@ -259,14 +261,6 @@ void __kmpc_push_target_tripcount(int64_t device_id, uint64_t loop_tripcount); } #endif -#ifdef OMPTARGET_DEBUG -#include -#define DEBUGP(prefix, ...) \ - { \ - fprintf(stderr, "%s --> ", prefix); \ - fprintf(stderr, __VA_ARGS__); \ - } - #ifndef __STDC_FORMAT_MACROS #define __STDC_FORMAT_MACROS #endif @@ -293,6 +287,14 @@ void __kmpc_push_target_tripcount(int64_t device_id, uint64_t loop_tripcount); * // 16 digits for 64bit * (uintptr_t) ptr); */ + +#ifdef OMPTARGET_DEBUG +#include +#define DEBUGP(prefix, ...) \ + { \ + fprintf(stderr, "%s --> ", prefix); \ + fprintf(stderr, __VA_ARGS__); \ + } #else #define DEBUGP(prefix, ...) \ {} diff --git a/openmp/libomptarget/src/device.cpp b/openmp/libomptarget/src/device.cpp index 5753543..867083f 100644 --- a/openmp/libomptarget/src/device.cpp +++ b/openmp/libomptarget/src/device.cpp @@ -160,8 +160,10 @@ 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) { + 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,40 @@ 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 (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 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 309785a..ebec76c 100644 --- a/openmp/libomptarget/src/device.h +++ b/openmp/libomptarget/src/device.h @@ -177,8 +177,9 @@ 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 = true, - bool HasCloseModifier = false); + bool &IsNew, bool &IsHostPtr, bool IsImplicit, + bool UpdateRefCount, bool HasCloseModifier, + bool HasPresentModifier); 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 6b4549b..47971b9 100644 --- a/openmp/libomptarget/src/omptarget.cpp +++ b/openmp/libomptarget/src/omptarget.cpp @@ -308,6 +308,7 @@ 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 @@ -316,13 +317,26 @@ 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. - Pointer_TgtPtrBegin = Device.getOrAllocTgtPtr(HstPtrBase, HstPtrBase, - sizeof(void *), Pointer_IsNew, IsHostPtr, IsImplicit, UpdateRef, - HasCloseModifier); + // 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); 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" @@ -334,13 +348,15 @@ 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); - 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 @@ -459,13 +475,27 @@ 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); - 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) { + // 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")); + } bool DelEntry = IsLast || ForceDelete; diff --git a/openmp/libomptarget/src/private.h b/openmp/libomptarget/src/private.h index cb20d8c..7772175 100644 --- a/openmp/libomptarget/src/private.h +++ b/openmp/libomptarget/src/private.h @@ -80,9 +80,19 @@ typedef int (*TargetDataFuncPtrTy)(DeviceTy &, int32_t, void **, void **, int64_t *, int64_t *, void **, __tgt_async_info *); //////////////////////////////////////////////////////////////////////////////// -// implementation for fatal messages +// implementation for 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 new file mode 100644 index 0000000..1d61dc0 --- /dev/null +++ b/openmp/libomptarget/test/mapping/present/target.c @@ -0,0 +1,42 @@ +// 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 new file mode 100644 index 0000000..fd3107d --- /dev/null +++ b/openmp/libomptarget/test/mapping/present/target_data.c @@ -0,0 +1,42 @@ +// 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 new file mode 100644 index 0000000..d96e7a4 --- /dev/null +++ b/openmp/libomptarget/test/mapping/present/target_enter_data.c @@ -0,0 +1,41 @@ +// 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 new file mode 100644 index 0000000..86b7ad8 --- /dev/null +++ b/openmp/libomptarget/test/mapping/present/target_exit_data.c @@ -0,0 +1,40 @@ +// 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 new file mode 100644 index 0000000..22d8746 --- /dev/null +++ b/openmp/libomptarget/test/mapping/present/unified_shared_memory.c @@ -0,0 +1,41 @@ +// 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 new file mode 100644 index 0000000..5488888 --- /dev/null +++ b/openmp/libomptarget/test/mapping/present/zero_length_array_section.c @@ -0,0 +1,45 @@ +// 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 new file mode 100644 index 0000000..bedc6a2 --- /dev/null +++ b/openmp/libomptarget/test/mapping/present/zero_length_array_section_exit.c @@ -0,0 +1,43 @@ +// 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; +}