[Openmp-commits] [openmp] 5ab4398 - [OpenMP] Fix `omp target update` for array extension
Joel E. Denny via Openmp-commits
openmp-commits at lists.llvm.org
Wed Aug 5 07:04:44 PDT 2020
Author: Joel E. Denny
Date: 2020-08-05T10:03:31-04:00
New Revision: 5ab43989c353a2378910d20c7b88e44ea92b3aee
URL: https://github.com/llvm/llvm-project/commit/5ab43989c353a2378910d20c7b88e44ea92b3aee
DIFF: https://github.com/llvm/llvm-project/commit/5ab43989c353a2378910d20c7b88e44ea92b3aee.diff
LOG: [OpenMP] Fix `omp target update` for array extension
OpenMP TR8 sec. 2.15.6 "target update Construct", p. 183, L3-4 states:
> If the corresponding list item is not present in the device data
> environment and there is no present modifier in the clause, then no
> assignment occurs to or from the original list item.
L10-11 states:
> If a present modifier appears in the clause and the corresponding
> list item is not present in the device data environment then an
> error occurs and the program termintates.
(OpenMP 5.0 also has the first passage but without mention of the
present modifier of course.)
In both passages, I assume "is not present" includes the case of
partially but not entirely present. However, without this patch, the
target update directive misbehaves in this case both with and without
the present modifier. For example:
```
#pragma omp target enter data map(to:arr[0:3])
#pragma omp target update to(arr[0:5]) // might fail on data transfer
#pragma omp target update to(present:arr[0:5]) // might fail on data transfer
```
The problem is that `DeviceTy::getTgtPtrBegin` does not return a null
pointer in that case, so `target_data_update` sees the data as fully
present, and the data transfer then might fail depending on the target
device. However, without the present modifier, there should never be
a failure. Moreover, with the present modifier, there should always
be a failure, and the diagnostic should mention the present modifier.
This patch fixes `DeviceTy::getTgtPtrBegin` to return null when
`target_data_update` is the caller. I'm wondering if it should do the
same for more callers.
Reviewed By: grokos, jdoerfert
Differential Revision: https://reviews.llvm.org/D85246
Added:
openmp/libomptarget/test/mapping/present/target_update_array_extension.c
openmp/libomptarget/test/mapping/target_update_array_extension.c
Modified:
openmp/libomptarget/src/device.cpp
openmp/libomptarget/src/device.h
openmp/libomptarget/src/omptarget.cpp
Removed:
################################################################################
diff --git a/openmp/libomptarget/src/device.cpp b/openmp/libomptarget/src/device.cpp
index 55d2e0162c9f..575eff15333f 100644
--- a/openmp/libomptarget/src/device.cpp
+++ b/openmp/libomptarget/src/device.cpp
@@ -236,14 +236,16 @@ void *DeviceTy::getOrAllocTgtPtr(void *HstPtrBegin, void *HstPtrBase,
// 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/device.h b/openmp/libomptarget/src/device.h
index 2cbb2f83e365..655cf96d8f6f 100644
--- a/openmp/libomptarget/src/device.h
+++ b/openmp/libomptarget/src/device.h
@@ -182,7 +182,8 @@ struct DeviceTy {
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/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp
index 541fe0d52fbb..5f8a3a679e8a 100644
--- a/openmp/libomptarget/src/omptarget.cpp
+++ b/openmp/libomptarget/src/omptarget.cpp
@@ -670,8 +670,8 @@ int target_data_update(DeviceTy &Device, int32_t arg_num,
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
index 000000000000..3e90c40c15b6
--- /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 <stdio.h>
+
+#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
index 000000000000..f5748ce90e95
--- /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 <stdio.h>
+
+#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;
+}
More information about the Openmp-commits
mailing list