[OpenMP] Fix always,from and delete for data absent at exit

Without this patch, there's a runtime error for those map types at
exit from an "omp target data" or at "omp target exit data", but the
spec says the list item should be ignored.

This patch tests that fix in data_absent_at_exit.c, and it also
improves other testing for data that is not fully present at exit.

Reviewed By: grokos, RaviNarayanaswamy

Differential Revision: https://reviews.llvm.org/D96999

Author: jdenny-ornl

openmp/libomptarget/src/omptarget.cpp

@@ -545,6 +545,13 @@ int targetDataEnd(ident_t *loc, DeviceTy &Device, int32_t ArgNum,
       DP("Mapping does not exist (%s)\n",
          (HasPresentModifier ? "'present' map type modifier" : "ignored"));
       if (HasPresentModifier) {
+        // OpenMP 5.1, sec. 2.21.7.1 "map Clause", p. 350 L10-13:
+        // "If a map clause appears on a target, target data, target enter data
+        // or target exit data construct with a present map-type-modifier then
+        // on entry to the region if the corresponding list item does not appear
+        // in the device data environment then an error occurs and the program
+        // terminates."
+        //
         // This should be an error upon entering an "omp target exit data".  It
         // should not be an error upon exiting an "omp target data" or "omp
         // target".  For "omp target data", Clang thus doesn't include present
@@ -564,6 +571,14 @@ int targetDataEnd(ident_t *loc, DeviceTy &Device, int32_t ArgNum,
          DataSize, DPxPTR(TgtPtrBegin), (IsLast ? "" : " not"));
     }
 
+    // OpenMP 5.1, sec. 2.21.7.1 "map Clause", p. 351 L14-16:
+    // "If the map clause appears on a target, target data, or target exit data
+    // construct and a corresponding list item of the original list item is not
+    // present in the device data environment on exit from the region then the
+    // list item is ignored."
+    if (!TgtPtrBegin)
+      continue;
+
     bool DelEntry = IsLast || ForceDelete;
 
     if ((ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF) &&

openmp/libomptarget/test/mapping/data_absent_at_exit.c

@@ -0,0 +1,33 @@
+// RUN: %libomptarget-compile-run-and-check-aarch64-unknown-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-powerpc64-ibm-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-powerpc64le-ibm-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-x86_64-pc-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-nvptx64-nvidia-cuda
+
+#include <stdio.h>
+
+// OpenMP 5.1, sec. 2.21.7.1 "map Clause", p. 351 L14-16:
+// "If the map clause appears on a target, target data, or target exit data
+// construct and a corresponding list item of the original list item is not
+// present in the device data environment on exit from the region then the
+// list item is ignored."
+
+int main(void) {
+  int f = 5, r = 6, d = 7, af = 8;
+
+  // Check exit from omp target data.
+  // CHECK: f = 5, af = 8
+  #pragma omp target data map(from: f) map(always, from: af)
+  {
+    #pragma omp target exit data map(delete: f, af)
+  }
+  printf("f = %d, af = %d\n", f, af);
+
+  // Check omp target exit data.
+  // CHECK: f = 5, r = 6, d = 7, af = 8
+  #pragma omp target exit data map(from: f) map(release: r) map(delete: d) \
+                               map(always, from: af)
+  printf("f = %d, r = %d, d = %d, af = %d\n", f, r, d, af);
+
+  return 0;
+}

openmp/libomptarget/test/mapping/present/target_exit_data.c renamed to openmp/libomptarget/test/mapping/present/target_exit_data_delete.c

@@ -23,18 +23,18 @@ int main() {
   fprintf(stderr, "addr=%p, size=%ld\n", &i, sizeof i);
 
   // CHECK-NOT: Libomptarget
-#pragma omp target enter data map(alloc: i)
-#pragma omp target exit data map(present, release: i)
+  #pragma omp target enter data map(alloc: i)
+  #pragma omp target exit data map(present, delete: i)
 
-  // CHECK: i is present
-  fprintf(stderr, "i is present\n");
+  // CHECK: i was present
+  fprintf(stderr, "i was present\n");
 
   // CHECK: Libomptarget message: device mapping required by 'present' map type 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 exit data map(present, release: i)
+  #pragma omp target exit data map(present, delete: i)
 
-  // CHECK-NOT: i is present
-  fprintf(stderr, "i is present\n");
+  // CHECK-NOT: i was present
+  fprintf(stderr, "i was present\n");
 
   return 0;
 }

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

@@ -0,0 +1,40 @@
+// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-fail-aarch64-unknown-linux-gnu 2>&1 \
+// RUN: | %fcheck-aarch64-unknown-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-fail-powerpc64-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-fail-powerpc64le-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64le-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-x86_64-pc-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-fail-x86_64-pc-linux-gnu 2>&1 \
+// RUN: | %fcheck-x86_64-pc-linux-gnu
+
+#include <stdio.h>
+
+int main() {
+  int i;
+
+  // CHECK: addr=0x[[#%x,HOST_ADDR:]], size=[[#%u,SIZE:]]
+  fprintf(stderr, "addr=%p, size=%ld\n", &i, sizeof i);
+
+  // CHECK-NOT: Libomptarget
+  #pragma omp target enter data map(alloc: i)
+  #pragma omp target exit data map(present, release: i)
+
+  // CHECK: i was present
+  fprintf(stderr, "i was present\n");
+
+  // CHECK: Libomptarget message: device mapping required by 'present' map type 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 exit data map(present, release: i)
+
+  // CHECK-NOT: i was present
+  fprintf(stderr, "i was present\n");
+
+  return 0;
+}

openmp/libomptarget/test/mapping/target_data_array_extension_at_exit.c

@@ -19,8 +19,8 @@
 
 // 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
+// RUN: %libomptarget-run-x86_64-pc-linux-gnu 2>&1 \
+// RUN: | %fcheck-x86_64-pc-linux-gnu
 
 // --------------------------------------------------
 // Check extends after