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

[jdenny]

Summary

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.

[jdenny]

openmp/libomptarget/test/mapping/target_data_array_extension_at_exit.c

23

This was an accident in my D85342.

[grokos]

Is this the missing behavior from D83061? Looks good, thanks!

[RaviNarayanaswamy]

In D96999#2573081, @grokos wrote:

Is this the missing behavior from D83061? Looks good, thanks!

Can you print a message under debug stating pointer does not exist

openmp/libomptarget/src/omptarget.cpp

579–580

Can you print a message under debug pointer does not exist.

[jdenny]

In D96999#2573081, @grokos wrote:

Is this the missing behavior from D83061?

You mean the map clause reordering stuff? No, as far as I know, it has no effect on that. I believe @ABataev was working on something along those lines, but unfortunately I lost track.

Looks good, thanks!

Thanks!

openmp/libomptarget/src/omptarget.cpp

579–580

I'm happy to add an additional message, but what do you want beyond the DP invocations in the preceding if and else?

[RaviNarayanaswamy]

openmp/libomptarget/src/omptarget.cpp

579–580

The previous DP should be fine. Thanks

[jdenny]

openmp/libomptarget/src/omptarget.cpp

579–580

OK, either way works for me. If you have a specific string, I'd be happy to add it before I push tomorrow.

openmp/libomptarget/src/omptarget.cpp

for (int32_t I = ArgNum - 1; I >= 0; --I) {
 
     // If PTR_AND_OBJ, HstPtrBegin is address of pointee
     void *TgtPtrBegin = Device.getTgtPtrBegin(
         HstPtrBegin, DataSize, IsLast, UpdateRef, IsHostPtr, !IsImplicit);
     if (!TgtPtrBegin && (DataSize || HasPresentModifier)) {
       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
         // modifiers for end calls.  For "omp target", we have not found a valid
         // OpenMP program for which the error matters: it appears that, if a
         // program can guarantee that data is present at the beginning of an
         // "omp target" region so that there's no error there, that data is also
         // guaranteed to be present at the end.
         MESSAGE("device mapping required by 'present' map type modifier does "
                 "not exist for host address " DPxMOD " (%" PRId64 " bytes)",
                 DPxPTR(HstPtrBegin), DataSize);
         return OFFLOAD_FAIL;
       }
     } else {
       DP("There are %" PRId64 " bytes allocated at target address " DPxMOD
          " - is%s last\n",
          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) &&
         !(ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ)) {
       DelEntry = false; // protect parent struct from being deallocated
     }
 
     if ((ArgTypes[I] & OMP_TGT_MAPTYPE_FROM) || DelEntry) {

openmp/libomptarget/test/mapping/data_absent_at_exit.c

+// 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_delete.c

 
 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)
+  #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

 
 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 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)
 
-  // 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/target_data_array_extension_at_exit.c

 
 // 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
+// RUN: %libomptarget-run-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-aarch64-unknown-linux-gnu 2>&1 \