[Openmp-commits] [openmp] 2cb926a - [OpenMP] Implement TR8 `present` motion modifier in runtime (2/2)

[Joel E. Denny via Openmp-commits]

Author: Joel E. Denny
Date: 2020-07-28T19:15:18-04:00
New Revision: 2cb926a447d22166dac0d6e6dceaf5b3616ef6a1

URL: https://github.com/llvm/llvm-project/commit/2cb926a447d22166dac0d6e6dceaf5b3616ef6a1
DIFF: https://github.com/llvm/llvm-project/commit/2cb926a447d22166dac0d6e6dceaf5b3616ef6a1.diff

LOG: [OpenMP] Implement TR8 `present` motion modifier in runtime (2/2)

This patch implements OpenMP runtime support for the OpenMP TR8
`present` motion modifier for `omp target update` directives.  The
previous patch in this series implements Clang front end support.

Reviewed By: grokos

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

Added: 
    openmp/libomptarget/test/mapping/present/target_update.c

Modified: 
    openmp/libomptarget/src/omptarget.cpp

Removed: 
    


################################################################################
diff  --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp
index 25722d75f6ea..aa2bebdebd80 100644
--- a/openmp/libomptarget/src/omptarget.cpp
+++ b/openmp/libomptarget/src/omptarget.cpp
@@ -627,6 +627,12 @@ int target_data_update(DeviceTy &Device, int32_t arg_num,
         false, IsHostPtr);
     if (!TgtPtrBegin) {
       DP("hst data:" DPxMOD " not found, becomes a noop\n", DPxPTR(HstPtrBegin));
+      if (arg_types[i] & OMP_TGT_MAPTYPE_PRESENT) {
+        MESSAGE("device mapping required by 'present' motion modifier does not "
+                "exist for host address " DPxMOD " (%ld bytes)",
+                DPxPTR(HstPtrBegin), MapSize);
+        return OFFLOAD_FAIL;
+      }
       continue;
     }
 

diff  --git a/openmp/libomptarget/test/mapping/present/target_update.c b/openmp/libomptarget/test/mapping/present/target_update.c
new file mode 100644
index 000000000000..3a3575a27d01
--- /dev/null
+++ b/openmp/libomptarget/test/mapping/present/target_update.c
@@ -0,0 +1,73 @@
+// --------------------------------------------------
+// Check 'to'
+// --------------------------------------------------
+
+// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu \
+// RUN:   -fopenmp-version=51 -DCLAUSE=to
+// 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 -DCLAUSE=to
+// 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 -DCLAUSE=to
+// 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 -DCLAUSE=to
+// RUN: %libomptarget-run-fail-x86_64-pc-linux-gnu 2>&1 \
+// RUN: | %fcheck-x86_64-pc-linux-gnu
+
+// --------------------------------------------------
+// Check 'from'
+// --------------------------------------------------
+
+// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu \
+// RUN:   -fopenmp-version=51 -DCLAUSE=from
+// 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 -DCLAUSE=from
+// 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 -DCLAUSE=from
+// 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 -DCLAUSE=from
+// 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 update CLAUSE(present: i)
+#pragma omp target exit data map(delete: i)
+
+  // CHECK: i is present
+  fprintf(stderr, "i is present\n");
+
+  // CHECK: Libomptarget message: device mapping required by 'present' motion 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 update CLAUSE(present: i)
+
+  // CHECK-NOT: i is present
+  fprintf(stderr, "i is present\n");
+
+  return 0;
+}