[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

diff --git a/openmp/libomptarget/src/device.cpp b/openmp/libomptarget/src/device.cpp
index 7049463..5a01257 100644 (file)
--- 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 (file)
index 0000000..870be39
--- /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 (file)
index 0000000..3aef777
--- /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;
+}