[Openmp-commits] [openmp] 8c8bb12 - [OpenMP] Fix `target data` exit for array extension
Joel E. Denny via Openmp-commits
openmp-commits at lists.llvm.org
Wed Aug 5 13:51:50 PDT 2020
Author: Joel E. Denny
Date: 2020-08-05T16:51:25-04:00
New Revision: 8c8bb128dfd09f84b27b9e732cf1355582b1d019
URL: https://github.com/llvm/llvm-project/commit/8c8bb128dfd09f84b27b9e732cf1355582b1d019
DIFF: https://github.com/llvm/llvm-project/commit/8c8bb128dfd09f84b27b9e732cf1355582b1d019.diff
LOG: [OpenMP] Fix `target data` exit for array extension
For example:
```
#pragma omp target data map(tofrom:arr[0:100])
{
#pragma omp target exit data map(delete:arr[0:100])
#pragma omp target enter data map(alloc:arr[98:2])
}
```
Without this patch, the transfer at the end of the target data region
is broken and fails depending on the target device. According to my
read of the spec, the transfer shouldn't even be attempted because
`arr[0:100]` isn't (fully) present there. To fix that, this patch
makes `DeviceTy::getTgtPtrBegin` return null for this case.
Reviewed By: grokos
Differential Revision: https://reviews.llvm.org/D85342
Added:
openmp/libomptarget/test/mapping/target_data_array_extension_at_exit.c
Modified:
openmp/libomptarget/src/omptarget.cpp
Removed:
################################################################################
diff --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp
index f7805831fd1a..2b158ff34268 100644
--- a/openmp/libomptarget/src/omptarget.cpp
+++ b/openmp/libomptarget/src/omptarget.cpp
@@ -505,7 +505,8 @@ int targetDataEnd(DeviceTy &Device, int32_t ArgNum, void **ArgBases,
// If PTR_AND_OBJ, HstPtrBegin is address of pointee
void *TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBegin, DataSize, IsLast,
- UpdateRef, IsHostPtr);
+ UpdateRef, IsHostPtr,
+ /*MustContain=*/true);
if (!TgtPtrBegin && (DataSize || HasPresentModifier)) {
DP("Mapping does not exist (%s)\n",
(HasPresentModifier ? "'present' map type modifier" : "ignored"));
diff --git a/openmp/libomptarget/test/mapping/target_data_array_extension_at_exit.c b/openmp/libomptarget/test/mapping/target_data_array_extension_at_exit.c
new file mode 100644
index 000000000000..89200ee0cebb
--- /dev/null
+++ b/openmp/libomptarget/test/mapping/target_data_array_extension_at_exit.c
@@ -0,0 +1,140 @@
+// --------------------------------------------------
+// Check extends before
+// --------------------------------------------------
+
+// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu \
+// RUN: -fopenmp-version=51 -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: -fopenmp-version=51 -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: -fopenmp-version=51 -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: -fopenmp-version=51 -DEXTENDS=BEFORE
+// XUN: %libomptarget-run-x86_64-pc-linux-gnu 2>&1 \
+// XUN: | %fcheck-x86_64-pc-linux-gnu
+
+// --------------------------------------------------
+// Check extends after
+// --------------------------------------------------
+
+// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu \
+// RUN: -fopenmp-version=51 -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: -fopenmp-version=51 -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: -fopenmp-version=51 -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: -fopenmp-version=51 -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
+
+#define SIZE 100
+
+#if EXTENDS == BEFORE
+# define SMALL_BEG (SIZE-2)
+# define SMALL_END SIZE
+# define LARGE_BEG 0
+# define LARGE_END SIZE
+#elif EXTENDS == AFTER
+# define SMALL_BEG 0
+# define SMALL_END 2
+# define LARGE_BEG 0
+# define LARGE_END SIZE
+#else
+# error EXTENDS undefined
+#endif
+
+#define SMALL SMALL_BEG:(SMALL_END-SMALL_BEG)
+#define LARGE LARGE_BEG:(LARGE_END-LARGE_BEG)
+
+void check_not_present() {
+ int arr[SIZE];
+
+ for (int i = 0; i < SIZE; ++i)
+ arr[i] = 99;
+
+ // CHECK-LABEL: checking not present
+ fprintf(stderr, "checking not present\n");
+
+ // arr[LARGE] isn't (fully) present at the end of the target data region, so
+ // the device-to-host transfer should not be performed, or it might fail.
+#pragma omp target data map(tofrom: arr[LARGE])
+ {
+#pragma omp target exit data map(delete: arr[LARGE])
+#pragma omp target enter data map(alloc: arr[SMALL])
+#pragma omp target map(alloc: arr[SMALL])
+ for (int i = SMALL_BEG; i < SMALL_END; ++i)
+ arr[i] = 88;
+ }
+
+ // CHECK-NOT: Libomptarget
+ // CHECK-NOT: error
+ for (int i = 0; i < SIZE; ++i) {
+ if (arr[i] != 99)
+ fprintf(stderr, "error: arr[%d]=%d\n", i, arr[i]);
+ }
+}
+
+void check_is_present() {
+ int arr[SIZE];
+
+ for (int i = 0; i < SIZE; ++i)
+ arr[i] = 99;
+
+ // CHECK-LABEL: checking is present
+ fprintf(stderr, "checking is present\n");
+
+ // arr[SMALL] is (fully) present at the end of the target data region, and the
+ // device-to-host transfer should be performed only for it even though more
+ // of the array is then present.
+#pragma omp target data map(tofrom: arr[SMALL])
+ {
+#pragma omp target exit data map(delete: arr[SMALL])
+#pragma omp target enter data map(alloc: arr[LARGE])
+#pragma omp target map(alloc: arr[LARGE])
+ for (int i = LARGE_BEG; i < LARGE_END; ++i)
+ arr[i] = 88;
+ }
+
+ // CHECK-NOT: Libomptarget
+ // CHECK-NOT: error
+ for (int i = 0; i < SIZE; ++i) {
+ if (SMALL_BEG <= i && i < SMALL_END) {
+ if (arr[i] != 88)
+ fprintf(stderr, "error: arr[%d]=%d\n", i, arr[i]);
+ } else if (arr[i] != 99) {
+ fprintf(stderr, "error: arr[%d]=%d\n", i, arr[i]);
+ }
+ }
+}
+
+int main() {
+ check_not_present();
+ check_is_present();
+ return 0;
+}
More information about the Openmp-commits
mailing list