[OpenMP] Fix `present` diagnostic for array extension

[jdenny]

Summary

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
  ;
}

[grokos]

openmp/libomptarget/src/device.cpp

194–202

On second thoughts, maybe this should also become a MESSAGE. It's a diagnostic which will help the user understand why the present-related message follows.

[vzakhari]

openmp/libomptarget/src/device.cpp

197

Can you please use PRId64 instead? '%ld' tends to break Windows builds.

[jdenny]

openmp/libomptarget/src/device.cpp

194–202

Are you suggesting it generally or only when HasPresentModifier?

197

Thanks for catching that. I'll fix this here and propose another patch to fix all the other places I made this mistake.

[grokos]

openmp/libomptarget/src/device.cpp

194–202

Generally.

197

I'm preparing a patch to replace all instances of %ld with PRId64 for int64_t variables throughout libomptarget. There are many more instances than introduced by your patches.

[jdenny]

openmp/libomptarget/src/device.cpp

197

Thanks. I'll just fix this one then.

[jdenny]

Add user diagnostic for mapping extension. Use PRId64 where appropriate.

[vzakhari]

openmp/libomptarget/src/device.cpp

197

Thank you!

[grokos]

LGTM. Thanks!

[jdenny]

Thanks for the quick reviews!

openmp/libomptarget/src/device.cpp

if (lr.Flags.IsContained ||
     DP("Mapping exists%s with HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD ", "
         "Size=%" PRId64 ",%s RefCount=%s\n", (IsImplicit ? " (implicit)" : ""),
         DPxPTR(HstPtrBegin), DPxPTR(tp), Size,
         (UpdateRefCount ? " updated" : ""),
         HT.isRefCountInf() ? "INF" : std::to_string(HT.getRefCount()).c_str());
     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
     // not privatized use host address. Any explicitly mapped variables also use
     // host address where correctness is not impeded. In all other cases maps
     // are respected.
     // In addition to the mapping rules above, the close map modifier forces the
     // mapping of the variable to the device.

openmp/libomptarget/test/mapping/present/target_array_extension.c

+// --------------------------------------------------
+// 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;
+}

openmp/libomptarget/test/mapping/present/target_data_array_extension.c

+// --------------------------------------------------
+// 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;
+}