openmp: Add support for the 'present' modifier

This implements support for the OpenMP 5.1 'present' modifier, which can be
used in map clauses in the 'target', 'target data', 'target data enter' and
'target data exit' constructs, and in the 'to' and 'from' clauses of the
'target update' construct.  It is also supported in defaultmap.

The modifier triggers a fatal runtime error if the data specified by the
clause is not already present on the target device.  It can also be combined
with 'always' in map clauses.

2023-06-06  Kwok Cheung Yeung  <kcy@codesourcery.com>
	    Tobias Burnus  <tobias@codesourcery.com>

gcc/c/
	* c-parser.cc (c_parser_omp_clause_defaultmap,
	c_parser_omp_clause_map): Parse 'present'.
	(c_parser_omp_clause_to, c_parser_omp_clause_from): Remove.
	(c_parser_omp_clause_from_to): New; parse to/from clauses with
	optional present modifer.
	(c_parser_omp_all_clauses): Update call.
	(c_parser_omp_target_data, c_parser_omp_target_enter_data,
	c_parser_omp_target_exit_data): Handle new map enum values
	for 'present' mapping.

gcc/cp/
	* parser.cc (cp_parser_omp_clause_defaultmap,
	cp_parser_omp_clause_map): Parse 'present'.
	(cp_parser_omp_clause_from_to): New; parse to/from
	clauses with optional 'present' modifier.
	(cp_parser_omp_all_clauses): Update call.
	(cp_parser_omp_target_data, cp_parser_omp_target_enter_data,
	cp_parser_omp_target_exit_data): Handle new enum value for
	'present' mapping.
	* semantics.cc (finish_omp_target): Likewise.

gcc/fortran/
	* dump-parse-tree.cc (show_omp_namelist): Display 'present' map
	modifier.
	(show_omp_clauses): Display 'present' motion modifier for 'to'
	and 'from' clauses.

	* gfortran.h (enum gfc_omp_map_op): Add entries with 'present'
	modifiers.
	(struct gfc_omp_namelist): Add 'present_modifer'.
	* openmp.cc (gfc_match_motion_var_list): New, handles optional
	'present' modifier for to/from clauses.
	(gfc_match_omp_clauses): Call it for to/from clauses; parse 'present'
	in defaultmap and map clauses.
	(resolve_omp_clauses): Allow 'present' modifiers on 'target',
	'target data', 'target enter' and 'target exit'	directives.
	* trans-openmp.cc (gfc_trans_omp_clauses): Apply 'present' modifiers
	to tree node for 'map', 'to' and 'from'	clauses.  Apply 'present' for
	defaultmap.

gcc/
	* gimplify.cc (omp_notice_variable): Apply GOVD_MAP_ALLOC_ONLY flag
	and defaultmap flags if the defaultmap has GOVD_MAP_FORCE_PRESENT flag
	set.
	(omp_get_attachment): Handle map clauses with 'present' modifier.
	(omp_group_base): Likewise.
	(gimplify_scan_omp_clauses): Reorder present maps to come first.
	Set GOVD flags for present defaultmaps.
	(gimplify_adjust_omp_clauses_1): Set map kind for present defaultmaps.
	* omp-low.cc (scan_sharing_clauses): Handle 'always, present' map
	clauses.
	(lower_omp_target): Handle map clauses with 'present' modifier.
	Handle 'to' and 'from' clauses with 'present'.
	* tree-core.h (enum omp_clause_defaultmap_kind): Add
	OMP_CLAUSE_DEFAULTMAP_PRESENT defaultmap kind.
	* tree-pretty-print.cc (dump_omp_clause): Handle 'map', 'to' and
	'from' clauses with 'present' modifier.  Handle present defaultmap.
	* tree.h (OMP_CLAUSE_MOTION_PRESENT): New #define.

include/
	* gomp-constants.h (GOMP_MAP_FLAG_SPECIAL_5): New.
	(GOMP_MAP_FLAG_FORCE): Redefine.
	(GOMP_MAP_FLAG_PRESENT, GOMP_MAP_FLAG_ALWAYS_PRESENT): New.
	(enum gomp_map_kind): Add map kinds with 'present' modifiers.
	(GOMP_MAP_COPY_TO_P, GOMP_MAP_COPY_FROM_P): Evaluate to true for
	map variants with 'present'
	(GOMP_MAP_ALWAYS_TO_P, GOMP_MAP_ALWAYS_FROM_P): Evaluate to true
	for map variants with 'always, present' modifiers.
	(GOMP_MAP_ALWAYS): Redefine.
	(GOMP_MAP_FORCE_P, GOMP_MAP_PRESENT_P): New.

libgomp/
	* libgomp.texi (OpenMP 5.1 Impl. status): Set 'present' support for
	defaultmap to 'Y', add 'Y' entry for 'present' on to/from/map clauses.
	* target.c (gomp_to_device_kind_p): Add map kinds with 'present'
	modifier.
	(gomp_map_vars_existing): Use new GOMP_MAP_FORCE_P macro.
	(gomp_map_vars_internal, gomp_update, gomp_target_rev):
	Emit runtime error if memory region not present.
	* testsuite/libgomp.c-c++-common/target-present-1.c: New test.
	* testsuite/libgomp.c-c++-common/target-present-2.c: New test.
	* testsuite/libgomp.c-c++-common/target-present-3.c: New test.
	* testsuite/libgomp.fortran/target-present-1.f90: New test.
	* testsuite/libgomp.fortran/target-present-2.f90: New test.
	* testsuite/libgomp.fortran/target-present-3.f90: New test.

gcc/testsuite/

	* c-c++-common/gomp/map-6.c: Update dg-error, extend to test for
	duplicated 'present' and extend scan-dump tests for 'present'.
	* gfortran.dg/gomp/defaultmap-1.f90: Update dg-error.
	* gfortran.dg/gomp/map-7.f90: Extend parse and dump test for
	'present'.
	* gfortran.dg/gomp/map-8.f90: Extend for duplicate 'present'
	modifier checking.
	* c-c++-common/gomp/defaultmap-4.c: New test.
	* c-c++-common/gomp/map-9.c: New test.
	* c-c++-common/gomp/target-update-1.c: New test.
	* gfortran.dg/gomp/defaultmap-8.f90: New test.
	* gfortran.dg/gomp/map-11.f90: New test.
	* gfortran.dg/gomp/map-12.f90: New test.
	* gfortran.dg/gomp/target-update-1.f90: New test.

8 files changed, 227 insertions, 6 deletions

diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi
index 3ea17a4..76c56a7 100644
--- a/libgomp/libgomp.texi
+++ b/libgomp/libgomp.texi
@@ -311,7 +311,7 @@ The OpenMP 4.5 specification is fully supported.
 @item @code{inoutset} argument to the @code{depend} clause @tab Y @tab
 @item @code{private} and @code{firstprivate} argument to @code{default}
       clause in C and C++ @tab Y @tab
-@item @code{present} argument to @code{defaultmap} clause @tab N @tab
+@item @code{present} argument to @code{defaultmap} clause @tab Y @tab
 @item @code{omp_set_num_teams}, @code{omp_set_teams_thread_limit},
       @code{omp_get_max_teams}, @code{omp_get_teams_thread_limit} runtime
       routines @tab Y @tab
@@ -353,6 +353,8 @@ to address of matching mapped list item per 5.1, Sect. 2.21.7.2 @tab N @tab
 @item Optional comma between directive and clause in the @code{#pragma} form @tab Y @tab
 @item @code{indirect} clause in @code{declare target} @tab N @tab
 @item @code{device_type(nohost)}/@code{device_type(host)} for variables @tab N @tab
+@item @code{present} modifier to the @code{map}, @code{to} and @code{from}
+      clauses @tab Y @tab
 @end multitable
 
 
diff --git a/libgomp/target.c b/libgomp/target.c
index 3238954..a9e8005 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -358,6 +358,8 @@ gomp_to_device_kind_p (int kind)
     case GOMP_MAP_FORCE_ALLOC:
     case GOMP_MAP_FORCE_FROM:
     case GOMP_MAP_ALWAYS_FROM:
+    case GOMP_MAP_PRESENT_FROM:
+    case GOMP_MAP_ALWAYS_PRESENT_FROM:
       return false;
     default:
       return true;
@@ -593,7 +595,7 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep,
   else
     tgt_var->length = newn->host_end - newn->host_start;
 
-  if ((kind & GOMP_MAP_FLAG_FORCE)
+  if (GOMP_MAP_FORCE_P (kind)
       /* For implicit maps, old contained in new is valid.  */
       || !(implicit_subset
 	   /* Otherwise, new contained inside old is considered valid.  */
@@ -1714,6 +1716,20 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
 #endif
 		    }
 		    break;
+		  case GOMP_MAP_PRESENT_ALLOC:
+		  case GOMP_MAP_PRESENT_TO:
+		  case GOMP_MAP_PRESENT_FROM:
+		  case GOMP_MAP_PRESENT_TOFROM:
+		  case GOMP_MAP_ALWAYS_PRESENT_TO:
+		  case GOMP_MAP_ALWAYS_PRESENT_FROM:
+		  case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
+		    /* We already looked up the memory region above and it
+		       was missing.  */
+		    gomp_mutex_unlock (&devicep->lock);
+		    gomp_fatal ("present clause: not present on the device "
+				"(%p, %d)",
+				(void *) k->host_start, devicep->target_id);
+		    break;
 		  case GOMP_MAP_FORCE_DEVICEPTR:
 		    assert (k->host_end - k->host_start == sizeof (void *));
 		    gomp_copy_host2dev (devicep, aq,
@@ -2124,6 +2140,20 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs,
 		  gomp_copy_dev2host (devicep, NULL, hostaddr, devaddr, size);
 	      }
 	  }
+	else
+	  {
+	    int kind = get_kind (short_mapkind, kinds, i);
+
+	    if (GOMP_MAP_PRESENT_P (kind))
+	      {
+		/* We already looked up the memory region above and it
+		   was missing.  */
+		gomp_mutex_unlock (&devicep->lock);
+		gomp_fatal ("present clause: not present on the device "
+			    "(%p, %d)",
+			    (void *) hostaddrs[i], devicep->target_id);
+	      }
+	  }
       }
   gomp_mutex_unlock (&devicep->lock);
 }
@@ -3422,7 +3452,8 @@ gomp_target_rev (uint64_t fn_ptr, uint64_t mapnum, uint64_t devaddrs_ptr,
 	      case GOMP_MAP_DELETE:
 	      case GOMP_MAP_RELEASE:
 	      case GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION:
-		/* Assume it is present; look it up - but ignore otherwise. */
+		/* Assume it is present; look it up - but ignore unless the
+		   present clause is there. */
 	      case GOMP_MAP_ALLOC:
 	      case GOMP_MAP_FROM:
 	      case GOMP_MAP_FORCE_ALLOC:
@@ -3434,6 +3465,12 @@ gomp_target_rev (uint64_t fn_ptr, uint64_t mapnum, uint64_t devaddrs_ptr,
 	      case GOMP_MAP_FORCE_TOFROM:
 	      case GOMP_MAP_ALWAYS_TO:
 	      case GOMP_MAP_ALWAYS_TOFROM:
+	      case GOMP_MAP_PRESENT_FROM:
+	      case GOMP_MAP_PRESENT_TO:
+	      case GOMP_MAP_PRESENT_TOFROM:
+	      case GOMP_MAP_ALWAYS_PRESENT_FROM:
+	      case GOMP_MAP_ALWAYS_PRESENT_TO:
+	      case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
 	      case GOMP_MAP_ZERO_LEN_ARRAY_SECTION:
 		cdata[i].devaddr = devaddrs[i];
 		bool zero_len = (kind == GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION
@@ -3454,7 +3491,23 @@ gomp_target_rev (uint64_t fn_ptr, uint64_t mapnum, uint64_t devaddrs_ptr,
 					      devaddrs[i] + sizes[i], zero_len);
 		    cdata[i].present = n2 != NULL;
 		  }
-		if (!cdata[i].present
+		if (!cdata[i].present && GOMP_MAP_PRESENT_P (kind))
+		  {
+		    gomp_mutex_unlock (&devicep->lock);
+#ifdef HAVE_INTTYPES_H
+		    gomp_fatal ("present clause: no corresponding data on "
+				"parent device at %p with size %"PRIu64,
+				(void *) (uintptr_t) devaddrs[i],
+				(uint64_t) sizes[i]);
+#else
+		    gomp_fatal ("present clause: no corresponding data on "
+				"parent device at %p with size %lu",
+				(void *) (uintptr_t) devaddrs[i],
+				(unsigned long) sizes[i]);
+#endif
+		    break;
+		  }
+		else if (!cdata[i].present
 		    && kind != GOMP_MAP_DELETE
 		    && kind != GOMP_MAP_RELEASE
 		    && kind != GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION)
@@ -3472,8 +3525,7 @@ gomp_target_rev (uint64_t fn_ptr, uint64_t mapnum, uint64_t devaddrs_ptr,
 		     && (kind == GOMP_MAP_TO || kind == GOMP_MAP_TOFROM))
 		    || kind == GOMP_MAP_FORCE_TO
 		    || kind == GOMP_MAP_FORCE_TOFROM
-		    || kind == GOMP_MAP_ALWAYS_TO
-		    || kind == GOMP_MAP_ALWAYS_TOFROM)
+		    || GOMP_MAP_ALWAYS_TO_P (kind))
 		  {
 		    gomp_copy_dev2host (devicep, aq,
 					(void *) (uintptr_t) devaddrs[i],
@@ -3658,6 +3710,10 @@ gomp_target_rev (uint64_t fn_ptr, uint64_t mapnum, uint64_t devaddrs_ptr,
 	      case GOMP_MAP_FORCE_TOFROM:
 	      case GOMP_MAP_ALWAYS_FROM:
 	      case GOMP_MAP_ALWAYS_TOFROM:
+	      case GOMP_MAP_PRESENT_FROM:
+	      case GOMP_MAP_PRESENT_TOFROM:
+	      case GOMP_MAP_ALWAYS_PRESENT_FROM:
+	      case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
 		copy = true;
 		/* FALLTHRU */
 	      case GOMP_MAP_FROM:
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-present-1.c b/libgomp/testsuite/libgomp.c-c++-common/target-present-1.c
new file mode 100644
index 0000000..12f154c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-present-1.c
@@ -0,0 +1,27 @@
+/* { dg-do run { target offload_device } } */
+/* { dg-shouldfail "present error triggered" } */
+
+#define N 100
+
+int main (void)
+{
+  int a[N], b[N], c[N];
+
+  for (int i = 0; i < N; i++) {
+    a[i] = i * 2;
+    b[i] = i * 3 + 1;
+  }
+
+  #pragma omp target enter data map (alloc: a, c)
+    /* a has already been allocated, so this should be okay.  */
+    #pragma omp target map (present, to: a)
+      for (int i = 0; i < N; i++)
+	c[i] = a[i];
+
+    /* b has not been allocated, so this should result in an error.  */
+    /* { dg-output "libgomp: present clause: not present on the device \\\(0x\[0-9a-f\]+, \[0-9\]+\\\)" } */
+    #pragma omp target map (present, to: b)
+      for (int i = 0; i < N; i++)
+	c[i] += b[i];
+  #pragma omp target exit data map (from: c)
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-present-2.c b/libgomp/testsuite/libgomp.c-c++-common/target-present-2.c
new file mode 100644
index 0000000..d4debba
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-present-2.c
@@ -0,0 +1,27 @@
+/* { dg-do run { target offload_device } } */
+/* { dg-shouldfail "present error triggered" } */
+
+#define N 100
+
+int main (void)
+{
+  int a[N], b[N], c[N];
+
+  for (int i = 0; i < N; i++) {
+    a[i] = i * 2;
+    b[i] = i * 3 + 1;
+  }
+
+  #pragma omp target enter data map (alloc: a, c)
+    /* a has already been allocated, so this should be okay.  */
+    #pragma omp target defaultmap (present)
+      for (int i = 0; i < N; i++)
+	c[i] = a[i];
+
+    /* b has not been allocated, so this should result in an error.  */
+    /* { dg-output "libgomp: present clause: not present on the device \\\(0x\[0-9a-f\]+, \[0-9\]+\\\)" } */
+    #pragma omp target defaultmap (present)
+      for (int i = 0; i < N; i++)
+	c[i] += b[i];
+  #pragma omp target exit data map (from: c)
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-present-3.c b/libgomp/testsuite/libgomp.c-c++-common/target-present-3.c
new file mode 100644
index 0000000..9d8d8f8
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-present-3.c
@@ -0,0 +1,27 @@
+/* { dg-do run { target offload_device } } */
+/* { dg-shouldfail "present error triggered" } */
+
+#include <stdio.h>
+
+#define N 100
+
+int main (void)
+{
+  int a[N], b[N], c[N];
+
+  for (int i = 0; i < N; i++) {
+    a[i] = i * 2;
+    b[i] = i * 3 + 1;
+  }
+
+  #pragma omp target enter data map (alloc: a, c)
+
+  /* This should work as a has already been allocated.  */
+  #pragma omp target update to (present: a)
+
+  /* This should fail as b has not been allocated.  */
+  /* { dg-output "libgomp: present clause: not present on the device \\\(0x\[0-9a-f\]+, \[0-9\]+\\\)" } */
+  #pragma omp target update to (present: b)
+
+  #pragma omp target exit data map (from: c)
+}
diff --git a/libgomp/testsuite/libgomp.fortran/target-present-1.f90 b/libgomp/testsuite/libgomp.fortran/target-present-1.f90
new file mode 100644
index 0000000..349dcb1
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-present-1.f90
@@ -0,0 +1,30 @@
+! { dg-do run { target offload_device } }
+! { dg-shouldfail "present error triggered" }
+
+program main
+  implicit none
+  integer, parameter :: N = 100
+  integer :: a(N), b(N), c(N), i
+
+  do i = 1, N
+    a(i) = i * 2
+    b(i) = i * 3 + 1
+  end do
+
+  !$omp target enter data map (alloc: a)
+    ! a has already been allocated, so this should be okay.
+    !$omp target map (present, to: a)
+      do i = 1, N
+        c(i) = a(i)
+      end do
+    !$omp end target
+
+    ! b has not been allocated, so this should result in an error.
+    ! { dg-output "libgomp: present clause: not present on the device \\\(0x\[0-9a-f\]+, \[0-9\]+\\\)" }
+    !$omp target map (present, to: b)
+      do i = 1, N
+        c(i) = c(i) + b(i)
+      end do
+    !$omp end target
+  !$omp target exit data map (from: c)
+end program
diff --git a/libgomp/testsuite/libgomp.fortran/target-present-2.f90 b/libgomp/testsuite/libgomp.fortran/target-present-2.f90
new file mode 100644
index 0000000..07e79d1
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-present-2.f90
@@ -0,0 +1,30 @@
+! { dg-do run { target offload_device } }
+! { dg-shouldfail "present error triggered" }
+
+program main
+  implicit none
+  integer, parameter :: N = 100
+  integer :: a(N), b(N), c(N), i
+
+  do i = 1, N
+    a(i) = i * 2
+    b(i) = i * 3 + 1
+  end do
+
+  !$omp target enter data map (alloc: a)
+    ! a has already been allocated, so this should be okay.
+    !$omp target defaultmap (present)
+      do i = 1, N
+        c(i) = a(i)
+      end do
+    !$omp end target
+
+    ! b has not been allocated, so this should result in an error.
+    ! { dg-output "libgomp: present clause: not present on the device \\\(0x\[0-9a-f\]+, \[0-9\]+\\\)" }
+    !$omp target defaultmap (present)
+      do i = 1, N
+        c(i) = c(i) + b(i)
+      end do
+    !$omp end target
+!$omp target exit data map (from: c)
+end program
diff --git a/libgomp/testsuite/libgomp.fortran/target-present-3.f90 b/libgomp/testsuite/libgomp.fortran/target-present-3.f90
new file mode 100644
index 0000000..a2709eb
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-present-3.f90
@@ -0,0 +1,22 @@
+! { dg-do run { target offload_device } }
+! { dg-shouldfail "present error triggered" }
+
+program main
+  implicit none
+  integer, parameter :: N = 100
+  integer :: a(N), b(N), c(N), i
+
+  do i = 1, N
+    a(i) = i * 2
+    b(i) = i * 3 + 1
+  end do
+
+  !$omp target enter data map (alloc: a, c)
+    ! This should work as a has already been allocated.
+    !$omp target update to (present: a)
+
+    ! This should fail as b has not been allocated.
+    ! { dg-output "libgomp: present clause: not present on the device \\\(0x\[0-9a-f\]+, \[0-9\]+\\\)" }
+    !$omp target update to (present: b)
+  !$omp target exit data map (from: c)
+end program