Index: openmp/libomptarget/src/omptarget.cpp =================================================================== --- openmp/libomptarget/src/omptarget.cpp +++ openmp/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) && Index: openmp/libomptarget/test/mapping/data_absent_at_exit.c =================================================================== --- /dev/null +++ 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 + +// 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; +} Index: openmp/libomptarget/test/mapping/present/target_exit_data_delete.c =================================================================== --- openmp/libomptarget/test/mapping/present/target_exit_data_delete.c +++ openmp/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; } Index: openmp/libomptarget/test/mapping/present/target_exit_data_release.c =================================================================== --- /dev/null +++ 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 + +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; +} Index: openmp/libomptarget/test/mapping/target_data_array_extension_at_exit.c =================================================================== --- openmp/libomptarget/test/mapping/target_data_array_extension_at_exit.c +++ 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