[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
GitOrigin-RevId: d2147b1a876187e6addd89b681d47f9f98a89669
diff --git a/libomptarget/src/omptarget.cpp b/libomptarget/src/omptarget.cpp
index d2d0e08..3d218ab 100644
--- a/libomptarget/src/omptarget.cpp
+++ b/libomptarget/src/omptarget.cpp
@@ -545,6 +545,13 @@
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 @@
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) &&
diff --git a/libomptarget/test/mapping/data_absent_at_exit.c b/libomptarget/test/mapping/data_absent_at_exit.c
new file mode 100644
index 0000000..a72bcfd
--- /dev/null
+++ b/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;
+}
diff --git a/libomptarget/test/mapping/present/target_exit_data.c b/libomptarget/test/mapping/present/target_exit_data_delete.c
similarity index 80%
copy from libomptarget/test/mapping/present/target_exit_data.c
copy to libomptarget/test/mapping/present/target_exit_data_delete.c
index 86b7ad8..761ee51 100644
--- a/libomptarget/test/mapping/present/target_exit_data.c
+++ b/libomptarget/test/mapping/present/target_exit_data_delete.c
@@ -23,18 +23,18 @@
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;
}
diff --git a/libomptarget/test/mapping/present/target_exit_data.c b/libomptarget/test/mapping/present/target_exit_data_release.c
similarity index 80%
rename from libomptarget/test/mapping/present/target_exit_data.c
rename to libomptarget/test/mapping/present/target_exit_data_release.c
index 86b7ad8..1f17732 100644
--- a/libomptarget/test/mapping/present/target_exit_data.c
+++ b/libomptarget/test/mapping/present/target_exit_data_release.c
@@ -23,18 +23,18 @@
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, 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)
+ #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;
}
diff --git a/libomptarget/test/mapping/target_data_array_extension_at_exit.c b/libomptarget/test/mapping/target_data_array_extension_at_exit.c
index 89200ee..d95a6b5 100644
--- a/libomptarget/test/mapping/target_data_array_extension_at_exit.c
+++ b/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
