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.

1 files changed, 105 insertions, 18 deletions

diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 5baa501..72f6fba 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -14939,6 +14939,13 @@ c_parser_omp_clause_defaultmap (c_parser *parser, tree list)
 	goto invalid_behavior;
       break;
 
+    case 'p':
+      if (strcmp ("present", p) == 0)
+	behavior = OMP_CLAUSE_DEFAULTMAP_PRESENT;
+      else
+	goto invalid_behavior;
+      break;
+
     case 't':
       if (strcmp ("tofrom", p) == 0)
 	behavior = OMP_CLAUSE_DEFAULTMAP_TOFROM;
@@ -17109,6 +17116,7 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
 
   int always_modifier = 0;
   int close_modifier = 0;
+  int present_modifier = 0;
   for (int pos = 1; pos < map_kind_pos; ++pos)
     {
       c_token *tok = c_parser_peek_token (parser);
@@ -17140,11 +17148,21 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
 	    }
 	  close_modifier++;
 	}
+      else if (strcmp ("present", p) == 0)
+	{
+	  if (present_modifier)
+	    {
+	      c_parser_error (parser, "too many %<present%> modifiers");
+	      parens.skip_until_found_close (parser);
+	      return list;
+	    }
+	  present_modifier++;
+	}
       else
 	{
 	  c_parser_error (parser, "%<#pragma omp target%> with "
-				  "modifier other than %<always%> or "
-				  "%<close%> on %<map%> clause");
+				  "modifier other than %<always%>, %<close%> "
+				  "or %<present%> on %<map%> clause");
 	  parens.skip_until_found_close (parser);
 	  return list;
 	}
@@ -17156,14 +17174,25 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
       && c_parser_peek_2nd_token (parser)->type == CPP_COLON)
     {
       const char *p = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value);
+      int always_present_modifier = always_modifier && present_modifier;
+
       if (strcmp ("alloc", p) == 0)
-	kind = GOMP_MAP_ALLOC;
+	kind = present_modifier ? GOMP_MAP_PRESENT_ALLOC : GOMP_MAP_ALLOC;
       else if (strcmp ("to", p) == 0)
-	kind = always_modifier ? GOMP_MAP_ALWAYS_TO : GOMP_MAP_TO;
+	kind = (always_present_modifier ? GOMP_MAP_ALWAYS_PRESENT_TO
+		: present_modifier ? GOMP_MAP_PRESENT_TO
+		: always_modifier ? GOMP_MAP_ALWAYS_TO
+		: GOMP_MAP_TO);
       else if (strcmp ("from", p) == 0)
-	kind = always_modifier ? GOMP_MAP_ALWAYS_FROM : GOMP_MAP_FROM;
+	kind = (always_present_modifier ? GOMP_MAP_ALWAYS_PRESENT_FROM
+		: present_modifier ? GOMP_MAP_PRESENT_FROM
+		: always_modifier ? GOMP_MAP_ALWAYS_FROM
+		: GOMP_MAP_FROM);
       else if (strcmp ("tofrom", p) == 0)
-	kind = always_modifier ? GOMP_MAP_ALWAYS_TOFROM : GOMP_MAP_TOFROM;
+	kind = (always_present_modifier ? GOMP_MAP_ALWAYS_PRESENT_TOFROM
+		: present_modifier ? GOMP_MAP_PRESENT_TOFROM
+		: always_modifier ? GOMP_MAP_ALWAYS_TOFROM
+		: GOMP_MAP_TOFROM);
       else if (strcmp ("release", p) == 0)
 	kind = GOMP_MAP_RELEASE;
       else if (strcmp ("delete", p) == 0)
@@ -17419,21 +17448,42 @@ c_parser_omp_clause_device_type (c_parser *parser, tree list)
 }
 
 /* OpenMP 4.0:
-   to ( variable-list ) */
+   from ( variable-list )
+   to ( variable-list )
+
+   OpenMP 5.1:
+   from ( [present :] variable-list )
+   to ( [present :] variable-list ) */
 
 static tree
-c_parser_omp_clause_to (c_parser *parser, tree list)
+c_parser_omp_clause_from_to (c_parser *parser, enum omp_clause_code kind,
+			     tree list)
 {
-  return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_TO, list, true);
-}
+  location_t loc = c_parser_peek_token (parser)->location;
+  matching_parens parens;
+  if (!parens.require_open (parser))
+    return list;
 
-/* OpenMP 4.0:
-   from ( variable-list ) */
+  bool present = false;
+  c_token *token = c_parser_peek_token (parser);
 
-static tree
-c_parser_omp_clause_from (c_parser *parser, tree list)
-{
-  return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_FROM, list, true);
+  if (token->type == CPP_NAME
+      && strcmp (IDENTIFIER_POINTER (token->value), "present") == 0
+      && c_parser_peek_2nd_token (parser)->type == CPP_COLON)
+    {
+      present = true;
+      c_parser_consume_token (parser);
+      c_parser_consume_token (parser);
+    }
+
+  tree nl = c_parser_omp_variable_list (parser, loc, kind, list);
+  parens.skip_until_found_close (parser);
+
+  if (present)
+    for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
+      OMP_CLAUSE_MOTION_PRESENT (c) = 1;
+
+  return nl;
 }
 
 /* OpenMP 4.0:
@@ -17940,11 +17990,13 @@ c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask,
 	      clauses = nl;
 	    }
 	  else
-	    clauses = c_parser_omp_clause_to (parser, clauses);
+	    clauses = c_parser_omp_clause_from_to (parser, OMP_CLAUSE_TO,
+						   clauses);
 	  c_name = "to";
 	  break;
 	case PRAGMA_OMP_CLAUSE_FROM:
-	  clauses = c_parser_omp_clause_from (parser, clauses);
+	  clauses = c_parser_omp_clause_from_to (parser, OMP_CLAUSE_FROM,
+						 clauses);
 	  c_name = "from";
 	  break;
 	case PRAGMA_OMP_CLAUSE_UNIFORM:
@@ -21768,11 +21820,18 @@ c_parser_omp_target_data (location_t loc, c_parser *parser, bool *if_p)
 	  {
 	  case GOMP_MAP_TO:
 	  case GOMP_MAP_ALWAYS_TO:
+	  case GOMP_MAP_PRESENT_TO:
+	  case GOMP_MAP_ALWAYS_PRESENT_TO:
 	  case GOMP_MAP_FROM:
 	  case GOMP_MAP_ALWAYS_FROM:
+	  case GOMP_MAP_PRESENT_FROM:
+	  case GOMP_MAP_ALWAYS_PRESENT_FROM:
 	  case GOMP_MAP_TOFROM:
 	  case GOMP_MAP_ALWAYS_TOFROM:
+	  case GOMP_MAP_PRESENT_TOFROM:
+	  case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
 	  case GOMP_MAP_ALLOC:
+	  case GOMP_MAP_PRESENT_ALLOC:
 	    map_seen = 3;
 	    break;
 	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
@@ -21918,7 +21977,10 @@ c_parser_omp_target_enter_data (location_t loc, c_parser *parser,
 	  {
 	  case GOMP_MAP_TO:
 	  case GOMP_MAP_ALWAYS_TO:
+	  case GOMP_MAP_PRESENT_TO:
+	  case GOMP_MAP_ALWAYS_PRESENT_TO:
 	  case GOMP_MAP_ALLOC:
+	  case GOMP_MAP_PRESENT_ALLOC:
 	    map_seen = 3;
 	    break;
 	  case GOMP_MAP_TOFROM:
@@ -21929,6 +21991,14 @@ c_parser_omp_target_enter_data (location_t loc, c_parser *parser,
 	    OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_ALWAYS_TO);
 	    map_seen = 3;
 	    break;
+	  case GOMP_MAP_PRESENT_TOFROM:
+	    OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_PRESENT_TO);
+	    map_seen = 3;
+	    break;
+	  case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
+	    OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_ALWAYS_PRESENT_TO);
+	    map_seen = 3;
+	    break;
 	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
 	  case GOMP_MAP_ALWAYS_POINTER:
 	  case GOMP_MAP_ATTACH_DETACH:
@@ -22016,6 +22086,8 @@ c_parser_omp_target_exit_data (location_t loc, c_parser *parser,
 	  {
 	  case GOMP_MAP_FROM:
 	  case GOMP_MAP_ALWAYS_FROM:
+	  case GOMP_MAP_PRESENT_FROM:
+	  case GOMP_MAP_ALWAYS_PRESENT_FROM:
 	  case GOMP_MAP_RELEASE:
 	  case GOMP_MAP_DELETE:
 	    map_seen = 3;
@@ -22028,6 +22100,14 @@ c_parser_omp_target_exit_data (location_t loc, c_parser *parser,
 	    OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_ALWAYS_FROM);
 	    map_seen = 3;
 	    break;
+	  case GOMP_MAP_PRESENT_TOFROM:
+	    OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_PRESENT_FROM);
+	    map_seen = 3;
+	    break;
+	  case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
+	    OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_ALWAYS_PRESENT_FROM);
+	    map_seen = 3;
+	    break;
 	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
 	  case GOMP_MAP_ALWAYS_POINTER:
 	  case GOMP_MAP_ATTACH_DETACH:
@@ -22273,11 +22353,18 @@ check_clauses:
 	  {
 	  case GOMP_MAP_TO:
 	  case GOMP_MAP_ALWAYS_TO:
+	  case GOMP_MAP_PRESENT_TO:
+	  case GOMP_MAP_ALWAYS_PRESENT_TO:
 	  case GOMP_MAP_FROM:
 	  case GOMP_MAP_ALWAYS_FROM:
+	  case GOMP_MAP_PRESENT_FROM:
+	  case GOMP_MAP_ALWAYS_PRESENT_FROM:
 	  case GOMP_MAP_TOFROM:
 	  case GOMP_MAP_ALWAYS_TOFROM:
+	  case GOMP_MAP_PRESENT_TOFROM:
+	  case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
 	  case GOMP_MAP_ALLOC:
+	  case GOMP_MAP_PRESENT_ALLOC:
 	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
 	  case GOMP_MAP_ALWAYS_POINTER:
 	  case GOMP_MAP_ATTACH_DETACH: