diff --git a/openmp/libomptarget/src/device.h b/openmp/libomptarget/src/device.h --- a/openmp/libomptarget/src/device.h +++ b/openmp/libomptarget/src/device.h @@ -182,7 +182,8 @@ bool HasPresentModifier); void *getTgtPtrBegin(void *HstPtrBegin, int64_t Size); void *getTgtPtrBegin(void *HstPtrBegin, int64_t Size, bool &IsLast, - bool UpdateRefCount, bool &IsHostPtr); + bool UpdateRefCount, bool &IsHostPtr, + bool MustContain = false); int deallocTgtPtr(void *TgtPtrBegin, int64_t Size, bool ForceDelete, bool HasCloseModifier = false); int associatePtr(void *HstPtrBegin, void *TgtPtrBegin, int64_t Size); 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 @@ -236,14 +236,16 @@ // Return the target pointer begin (where the data will be moved). // Decrement the reference counter if called from targetDataEnd. void *DeviceTy::getTgtPtrBegin(void *HstPtrBegin, int64_t Size, bool &IsLast, - bool UpdateRefCount, bool &IsHostPtr) { + bool UpdateRefCount, bool &IsHostPtr, + bool MustContain) { void *rc = NULL; IsHostPtr = false; IsLast = false; DataMapMtx.lock(); LookupResult lr = lookupMapping(HstPtrBegin, Size); - if (lr.Flags.IsContained || lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter) { + if (lr.Flags.IsContained || + (!MustContain && (lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter))) { auto &HT = *lr.Entry; IsLast = HT.getRefCount() == 1; 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 @@ -670,8 +670,8 @@ void *HstPtrBegin = args[i]; int64_t MapSize = arg_sizes[i]; bool IsLast, IsHostPtr; - void *TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBegin, MapSize, IsLast, - false, IsHostPtr); + void *TgtPtrBegin = Device.getTgtPtrBegin( + HstPtrBegin, MapSize, IsLast, false, IsHostPtr, /*MustContain=*/true); if (!TgtPtrBegin) { DP("hst data:" DPxMOD " not found, becomes a noop\n", DPxPTR(HstPtrBegin)); if (arg_types[i] & OMP_TGT_MAPTYPE_PRESENT) { diff --git a/openmp/libomptarget/test/mapping/present/target_update_array_extension.c b/openmp/libomptarget/test/mapping/present/target_update_array_extension.c new file mode 100644 --- /dev/null +++ b/openmp/libomptarget/test/mapping/present/target_update_array_extension.c @@ -0,0 +1,140 @@ +// -------------------------------------------------- +// Check 'to' and extends before +// -------------------------------------------------- + +// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu \ +// RUN: -fopenmp-version=51 -DCLAUSE=to -DEXTENDS=BEFORE +// RUN: %libomptarget-run-fail-aarch64-unknown-linux-gnu 2>&1 \ +// RUN: | %fcheck-aarch64-unknown-linux-gnu + +// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu \ +// RUN: -fopenmp-version=51 -DCLAUSE=to -DEXTENDS=BEFORE +// RUN: %libomptarget-run-fail-powerpc64-ibm-linux-gnu 2>&1 \ +// RUN: | %fcheck-powerpc64-ibm-linux-gnu + +// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu \ +// RUN: -fopenmp-version=51 -DCLAUSE=to -DEXTENDS=BEFORE +// 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: -fopenmp-version=51 -DCLAUSE=to -DEXTENDS=BEFORE +// RUN: %libomptarget-run-fail-x86_64-pc-linux-gnu 2>&1 \ +// RUN: | %fcheck-x86_64-pc-linux-gnu + +// -------------------------------------------------- +// Check 'from' and extends before +// -------------------------------------------------- + +// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu \ +// RUN: -fopenmp-version=51 -DCLAUSE=from -DEXTENDS=BEFORE +// RUN: %libomptarget-run-fail-aarch64-unknown-linux-gnu 2>&1 \ +// RUN: | %fcheck-aarch64-unknown-linux-gnu + +// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu \ +// RUN: -fopenmp-version=51 -DCLAUSE=from -DEXTENDS=BEFORE +// RUN: %libomptarget-run-fail-powerpc64-ibm-linux-gnu 2>&1 \ +// RUN: | %fcheck-powerpc64-ibm-linux-gnu + +// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu \ +// RUN: -fopenmp-version=51 -DCLAUSE=from -DEXTENDS=BEFORE +// 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: -fopenmp-version=51 -DCLAUSE=from -DEXTENDS=BEFORE +// RUN: %libomptarget-run-fail-x86_64-pc-linux-gnu 2>&1 \ +// RUN: | %fcheck-x86_64-pc-linux-gnu + +// -------------------------------------------------- +// Check 'to' and extends after +// -------------------------------------------------- + +// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu \ +// RUN: -fopenmp-version=51 -DCLAUSE=to -DEXTENDS=AFTER +// RUN: %libomptarget-run-fail-aarch64-unknown-linux-gnu 2>&1 \ +// RUN: | %fcheck-aarch64-unknown-linux-gnu + +// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu \ +// RUN: -fopenmp-version=51 -DCLAUSE=to -DEXTENDS=AFTER +// RUN: %libomptarget-run-fail-powerpc64-ibm-linux-gnu 2>&1 \ +// RUN: | %fcheck-powerpc64-ibm-linux-gnu + +// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu \ +// RUN: -fopenmp-version=51 -DCLAUSE=to -DEXTENDS=AFTER +// 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: -fopenmp-version=51 -DCLAUSE=to -DEXTENDS=AFTER +// RUN: %libomptarget-run-fail-x86_64-pc-linux-gnu 2>&1 \ +// RUN: | %fcheck-x86_64-pc-linux-gnu + +// -------------------------------------------------- +// Check 'from' and extends after +// -------------------------------------------------- + +// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu \ +// RUN: -fopenmp-version=51 -DCLAUSE=from -DEXTENDS=AFTER +// RUN: %libomptarget-run-fail-aarch64-unknown-linux-gnu 2>&1 \ +// RUN: | %fcheck-aarch64-unknown-linux-gnu + +// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu \ +// RUN: -fopenmp-version=51 -DCLAUSE=from -DEXTENDS=AFTER +// RUN: %libomptarget-run-fail-powerpc64-ibm-linux-gnu 2>&1 \ +// RUN: | %fcheck-powerpc64-ibm-linux-gnu + +// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu \ +// RUN: -fopenmp-version=51 -DCLAUSE=from -DEXTENDS=AFTER +// 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: -fopenmp-version=51 -DCLAUSE=from -DEXTENDS=AFTER +// RUN: %libomptarget-run-fail-x86_64-pc-linux-gnu 2>&1 \ +// RUN: | %fcheck-x86_64-pc-linux-gnu + +// END. + +#include + +#define BEFORE 0 +#define AFTER 1 + +#if EXTENDS == BEFORE +# define SMALL 2:3 +# define LARGE 0:5 +#elif EXTENDS == AFTER +# define SMALL 0:3 +# define LARGE 0:5 +#else +# error EXTENDS undefined +#endif + +int main() { + int arr[5]; + + // CHECK: addr=0x[[#%x,HOST_ADDR:]], size=[[#%u,SIZE:]] + fprintf(stderr, "addr=%p, size=%ld\n", arr, sizeof arr); + + // CHECK-NOT: Libomptarget +#pragma omp target data map(alloc: arr[LARGE]) + { +#pragma omp target update CLAUSE(present: arr[SMALL]) + } + + // CHECK: arr is present + fprintf(stderr, "arr is present\n"); + + // CHECK: Libomptarget message: device mapping required by 'present' motion 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(alloc: arr[SMALL]) + { +#pragma omp target update CLAUSE(present: arr[LARGE]) + } + + // CHECK-NOT: arr is present + fprintf(stderr, "arr is present\n"); + + return 0; +} diff --git a/openmp/libomptarget/test/mapping/target_update_array_extension.c b/openmp/libomptarget/test/mapping/target_update_array_extension.c new file mode 100644 --- /dev/null +++ b/openmp/libomptarget/test/mapping/target_update_array_extension.c @@ -0,0 +1,136 @@ +// -------------------------------------------------- +// Check 'to' and extends before +// -------------------------------------------------- + +// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu \ +// RUN: -DCLAUSE=to -DEXTENDS=BEFORE +// RUN: %libomptarget-run-aarch64-unknown-linux-gnu 2>&1 \ +// RUN: | %fcheck-aarch64-unknown-linux-gnu + +// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu \ +// RUN: -DCLAUSE=to -DEXTENDS=BEFORE +// RUN: %libomptarget-run-powerpc64-ibm-linux-gnu 2>&1 \ +// RUN: | %fcheck-powerpc64-ibm-linux-gnu + +// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu \ +// RUN: -DCLAUSE=to -DEXTENDS=BEFORE +// RUN: %libomptarget-run-powerpc64le-ibm-linux-gnu 2>&1 \ +// RUN: | %fcheck-powerpc64le-ibm-linux-gnu + +// RUN: %libomptarget-compile-x86_64-pc-linux-gnu \ +// RUN: -DCLAUSE=to -DEXTENDS=BEFORE +// RUN: %libomptarget-run-x86_64-pc-linux-gnu 2>&1 \ +// RUN: | %fcheck-x86_64-pc-linux-gnu + +// -------------------------------------------------- +// Check 'from' and extends before +// -------------------------------------------------- + +// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu \ +// RUN: -DCLAUSE=from -DEXTENDS=BEFORE +// RUN: %libomptarget-run-aarch64-unknown-linux-gnu 2>&1 \ +// RUN: | %fcheck-aarch64-unknown-linux-gnu + +// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu \ +// RUN: -DCLAUSE=from -DEXTENDS=BEFORE +// RUN: %libomptarget-run-powerpc64-ibm-linux-gnu 2>&1 \ +// RUN: | %fcheck-powerpc64-ibm-linux-gnu + +// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu \ +// RUN: -DCLAUSE=from -DEXTENDS=BEFORE +// RUN: %libomptarget-run-powerpc64le-ibm-linux-gnu 2>&1 \ +// RUN: | %fcheck-powerpc64le-ibm-linux-gnu + +// RUN: %libomptarget-compile-x86_64-pc-linux-gnu \ +// RUN: -DCLAUSE=from -DEXTENDS=BEFORE +// RUN: %libomptarget-run-x86_64-pc-linux-gnu 2>&1 \ +// RUN: | %fcheck-x86_64-pc-linux-gnu + +// -------------------------------------------------- +// Check 'to' and extends after +// -------------------------------------------------- + +// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu \ +// RUN: -DCLAUSE=to -DEXTENDS=AFTER +// RUN: %libomptarget-run-aarch64-unknown-linux-gnu 2>&1 \ +// RUN: | %fcheck-aarch64-unknown-linux-gnu + +// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu \ +// RUN: -DCLAUSE=to -DEXTENDS=AFTER +// RUN: %libomptarget-run-powerpc64-ibm-linux-gnu 2>&1 \ +// RUN: | %fcheck-powerpc64-ibm-linux-gnu + +// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu \ +// RUN: -DCLAUSE=to -DEXTENDS=AFTER +// RUN: %libomptarget-run-powerpc64le-ibm-linux-gnu 2>&1 \ +// RUN: | %fcheck-powerpc64le-ibm-linux-gnu + +// RUN: %libomptarget-compile-x86_64-pc-linux-gnu \ +// RUN: -DCLAUSE=to -DEXTENDS=AFTER +// RUN: %libomptarget-run-x86_64-pc-linux-gnu 2>&1 \ +// RUN: | %fcheck-x86_64-pc-linux-gnu + +// -------------------------------------------------- +// Check 'from' and extends after +// -------------------------------------------------- + +// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu \ +// RUN: -DCLAUSE=from -DEXTENDS=AFTER +// RUN: %libomptarget-run-aarch64-unknown-linux-gnu 2>&1 \ +// RUN: | %fcheck-aarch64-unknown-linux-gnu + +// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu \ +// RUN: -DCLAUSE=from -DEXTENDS=AFTER +// RUN: %libomptarget-run-powerpc64-ibm-linux-gnu 2>&1 \ +// RUN: | %fcheck-powerpc64-ibm-linux-gnu + +// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu \ +// RUN: -DCLAUSE=from -DEXTENDS=AFTER +// RUN: %libomptarget-run-powerpc64le-ibm-linux-gnu 2>&1 \ +// RUN: | %fcheck-powerpc64le-ibm-linux-gnu + +// RUN: %libomptarget-compile-x86_64-pc-linux-gnu \ +// RUN: -DCLAUSE=from -DEXTENDS=AFTER +// RUN: %libomptarget-run-x86_64-pc-linux-gnu 2>&1 \ +// RUN: | %fcheck-x86_64-pc-linux-gnu + +// END. + +#include + +#define BEFORE 0 +#define AFTER 1 + +#if EXTENDS == BEFORE +# define SMALL 2:3 +# define LARGE 0:5 +#elif EXTENDS == AFTER +# define SMALL 0:3 +# define LARGE 0:5 +#else +# error EXTENDS undefined +#endif + +int main() { + int arr[5]; + + // CHECK-NOT: Libomptarget +#pragma omp target data map(alloc: arr[LARGE]) + { +#pragma omp target update CLAUSE(arr[SMALL]) + } + + // CHECK: success + fprintf(stderr, "success\n"); + + // CHECK-NOT: Libomptarget +#pragma omp target data map(alloc: arr[SMALL]) + { +#pragma omp target update CLAUSE(arr[LARGE]) + } + + // CHECK: success + fprintf(stderr, "success\n"); + + return 0; +}