[Openmp-commits] [openmp] 41b1aef - [OpenMP] Fix `present` diagnostic for array extension
Joel E. Denny via Openmp-commits
openmp-commits at lists.llvm.org
Wed Aug 5 13:51:48 PDT 2020
Author: Joel E. Denny
Date: 2020-08-05T16:51:24-04:00
New Revision: 41b1aefecb9447620dd182b0352abed0df05665c
URL: https://github.com/llvm/llvm-project/commit/41b1aefecb9447620dd182b0352abed0df05665c
DIFF: https://github.com/llvm/llvm-project/commit/41b1aefecb9447620dd182b0352abed0df05665c.diff
LOG: [OpenMP] Fix `present` diagnostic for array extension
For example, without this patch, the following fails as expected with
or without the `present` modifier, but the `present` modifier doesn't
produce its usual diagnostic:
```
#pragma omp target data map(alloc: arr[0:2])
{
#pragma omp target map(present, tofrom: arr[0:100]) // not fully present
;
}
```
Reviewed By: grokos, vzakhari
Differential Revision: https://reviews.llvm.org/D85320
Added:
openmp/libomptarget/test/mapping/present/target_array_extension.c
openmp/libomptarget/test/mapping/present/target_data_array_extension.c
Modified:
openmp/libomptarget/src/device.cpp
Removed:
################################################################################
diff --git a/openmp/libomptarget/src/device.cpp b/openmp/libomptarget/src/device.cpp
index 70494637e29c..5a01257f132d 100644
--- a/openmp/libomptarget/src/device.cpp
+++ b/openmp/libomptarget/src/device.cpp
@@ -191,7 +191,15 @@ void *DeviceTy::getOrAllocTgtPtr(void *HstPtrBegin, void *HstPtrBase,
rc = (void *)tp;
} 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");
+ MESSAGE("explicit extension not allowed: host address specified is " DPxMOD
+ " (%" PRId64 " bytes), but device allocation maps to host at "
+ DPxMOD " (%" PRId64 " bytes)",
+ DPxPTR(HstPtrBegin), Size, DPxPTR(lr.Entry->HstPtrBegin),
+ lr.Entry->HstPtrEnd - lr.Entry->HstPtrBegin);
+ if (HasPresentModifier)
+ MESSAGE("device mapping required by 'present' map type modifier does not "
+ "exist for host address " DPxMOD " (%" PRId64 " bytes)",
+ DPxPTR(HstPtrBegin), Size);
} else if (RTLs->RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY &&
!HasCloseModifier) {
// If unified shared memory is active, implicitly mapped variables that are
diff --git a/openmp/libomptarget/test/mapping/present/target_array_extension.c b/openmp/libomptarget/test/mapping/present/target_array_extension.c
new file mode 100644
index 000000000000..870be3956c46
--- /dev/null
+++ b/openmp/libomptarget/test/mapping/present/target_array_extension.c
@@ -0,0 +1,112 @@
+// --------------------------------------------------
+// Check extends before
+// --------------------------------------------------
+
+// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu \
+// RUN: -fopenmp-version=51 -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 -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 -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 -DEXTENDS=BEFORE
+// RUN: %libomptarget-run-fail-x86_64-pc-linux-gnu 2>&1 \
+// RUN: | %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-fail-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-fail-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-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 -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
+
+#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_SIZE (SMALL_END-SMALL_BEG)
+#define LARGE_SIZE (LARGE_END-LARGE_BEG)
+
+#define SMALL SMALL_BEG:SMALL_SIZE
+#define LARGE LARGE_BEG:LARGE_SIZE
+
+int main() {
+ int arr[SIZE];
+
+ // CHECK: addr=0x[[#%x,SMALL_ADDR:]], size=[[#%u,SMALL_BYTES:]]
+ fprintf(stderr, "addr=%p, size=%ld\n", &arr[SMALL_BEG],
+ SMALL_SIZE * sizeof arr[0]);
+
+ // CHECK: addr=0x[[#%x,LARGE_ADDR:]], size=[[#%u,LARGE_BYTES:]]
+ fprintf(stderr, "addr=%p, size=%ld\n", &arr[LARGE_BEG],
+ LARGE_SIZE * sizeof arr[0]);
+
+ // CHECK-NOT: Libomptarget
+#pragma omp target data map(alloc: arr[LARGE])
+ {
+#pragma omp target map(present, tofrom: arr[SMALL])
+ ;
+ }
+
+ // CHECK: arr is present
+ fprintf(stderr, "arr is present\n");
+
+ // CHECK: Libomptarget message: explicit extension not allowed: host address specified is 0x{{0*}}[[#LARGE_ADDR]] ([[#LARGE_BYTES]] bytes), but device allocation maps to host at 0x{{0*}}[[#SMALL_ADDR]] ([[#SMALL_BYTES]] bytes)
+ // CHECK: Libomptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#LARGE_ADDR]] ([[#LARGE_BYTES]] 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 map(present, tofrom: arr[LARGE])
+ ;
+ }
+
+ // CHECK-NOT: arr is present
+ fprintf(stderr, "arr is present\n");
+
+ return 0;
+}
diff --git a/openmp/libomptarget/test/mapping/present/target_data_array_extension.c b/openmp/libomptarget/test/mapping/present/target_data_array_extension.c
new file mode 100644
index 000000000000..3aef7772f9e2
--- /dev/null
+++ b/openmp/libomptarget/test/mapping/present/target_data_array_extension.c
@@ -0,0 +1,112 @@
+// --------------------------------------------------
+// Check extends before
+// --------------------------------------------------
+
+// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu \
+// RUN: -fopenmp-version=51 -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 -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 -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 -DEXTENDS=BEFORE
+// RUN: %libomptarget-run-fail-x86_64-pc-linux-gnu 2>&1 \
+// RUN: | %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-fail-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-fail-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-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 -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
+
+#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_SIZE (SMALL_END-SMALL_BEG)
+#define LARGE_SIZE (LARGE_END-LARGE_BEG)
+
+#define SMALL SMALL_BEG:SMALL_SIZE
+#define LARGE LARGE_BEG:LARGE_SIZE
+
+int main() {
+ int arr[SIZE];
+
+ // CHECK: addr=0x[[#%x,SMALL_ADDR:]], size=[[#%u,SMALL_BYTES:]]
+ fprintf(stderr, "addr=%p, size=%ld\n", &arr[SMALL_BEG],
+ SMALL_SIZE * sizeof arr[0]);
+
+ // CHECK: addr=0x[[#%x,LARGE_ADDR:]], size=[[#%u,LARGE_BYTES:]]
+ fprintf(stderr, "addr=%p, size=%ld\n", &arr[LARGE_BEG],
+ LARGE_SIZE * sizeof arr[0]);
+
+ // CHECK-NOT: Libomptarget
+#pragma omp target data map(alloc: arr[LARGE])
+ {
+#pragma omp target data map(present, tofrom: arr[SMALL])
+ ;
+ }
+
+ // CHECK: arr is present
+ fprintf(stderr, "arr is present\n");
+
+ // CHECK: Libomptarget message: explicit extension not allowed: host address specified is 0x{{0*}}[[#LARGE_ADDR]] ([[#LARGE_BYTES]] bytes), but device allocation maps to host at 0x{{0*}}[[#SMALL_ADDR]] ([[#SMALL_BYTES]] bytes)
+ // CHECK: Libomptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#LARGE_ADDR]] ([[#LARGE_BYTES]] 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 data map(present, tofrom: arr[LARGE])
+ ;
+ }
+
+ // CHECK-NOT: arr is present
+ fprintf(stderr, "arr is present\n");
+
+ return 0;
+}
More information about the Openmp-commits
mailing list