OpenMP: Enable 'declare mapper' mappers for 'target update' directives

This patch enables use of 'declare mapper' for 'target update' directives,
for each of C, C++ and Fortran.

There are some implementation choices here and some
"read-between-the-lines" consequences regarding this functionality,
as follows:

 * It is possible to invoke a mapper which contains clauses that
   don't make sense for a given 'target update' operation.  E.g. if a
   mapper definition specifies a "from:" mapping and the user does "target
   update to(...)" which triggers that mapper, the resulting map kind
   (OpenMP 5.2, "Table 5.3: Map-Type Decay of Map Type Combinations")
   is "alloc" (and for the inverse case "release").  For such cases,
   an unconditional warning is issued and the map clause in question is
   dropped from the mapper expansion.  (Other choices might be to make
   this an error, or to do the same thing but silently, or warn only
   given some special option.)

 * The array-shaping operator is *permitted* for map clauses within
   'declare mapper' definitions.  That is because such mappers may be used
   for 'target update' directives, where the array-shaping operator is
   permitted.  I think that makes sense, depending on the semantic model
   of how and when substitution is supposed to take place, but I couldn't
   find such behaviour explicitly mentioned in the spec (as of 5.2).
   If the mapper is triggered by a different directive ("omp target",
   "omp target data", etc.), an error will be raised.

Support is also added for the "mapper" modifier on to/from clauses for
all three base languages.

2023-08-10  Julian Brown  <julian@codesourcery.com>

gcc/c-family/
	* c-common.h (c_omp_region_type): Add C_ORT_UPDATE and C_ORT_OMP_UPDATE
	codes.
	* c-omp.cc (omp_basic_map_kind_name): New function.
	(omp_instantiate_mapper): Add LOC parameter.  Add 'target update'
	support.
	(c_omp_instantiate_mappers): Add 'target update' support.

gcc/c/
	* c-parser.cc (c_parser_omp_variable_list): Support array-shaping
	operator in 'declare mapper' definitions.
	(c_parser_omp_clause_map): Pass C_ORT_OMP_DECLARE_MAPPER to
	c_parser_omp_variable_list in mapper definitions.
	(c_parser_omp_clause_from_to): Add parsing for mapper modifier.
	(c_parser_omp_target_update): Instantiate mappers.

gcc/cp/
	* parser.cc (cp_parser_omp_var_list_no_open): Support array-shaping
	operator in 'declare mapper' definitions.
	(cp_parser_omp_clause_from_to): Add parsing for mapper modifier.
	(cp_parser_omp_clause_map): Pass C_ORT_OMP_DECLARE_MAPPER to
	cp_parser_omp_var_list_no_open in mapper definitions.
	(cp_parser_omp_target_update): Instantiate mappers.

gcc/fortran/
	* openmp.cc (gfc_match_motion_var_list): Add parsing for mapper
	modifier.
	(gfc_match_omp_clauses): Adjust error handling for changes to
	gfc_match_motion_var_list.
	* trans-openmp.cc (gfc_trans_omp_clauses): Use correct ref for update
	operations.
	(gfc_trans_omp_target_update): Instantiate mappers.

gcc/testsuite/
	* c-c++-common/gomp/declare-mapper-17.c: New test.
	* c-c++-common/gomp/declare-mapper-19.c: New test.
	* gfortran.dg/gomp/declare-mapper-24.f90: New test.
	* gfortran.dg/gomp/declare-mapper-26.f90: Uncomment 'target update' part
	of test.
	* gfortran.dg/gomp/declare-mapper-27.f90: New test.

libgomp/
	* testsuite/libgomp.c-c++-common/declare-mapper-18.c: New test.
	* testsuite/libgomp.fortran/declare-mapper-25.f90: New test.
	* testsuite/libgomp.fortran/declare-mapper-28.f90: New test.

3 files changed, 116 insertions, 12 deletions

diff --git a/gcc/c-family/ChangeLog.omp b/gcc/c-family/ChangeLog.omp
index 07c30d0a9d1..60f5c29276f 100644
--- a/gcc/c-family/ChangeLog.omp
+++ b/gcc/c-family/ChangeLog.omp
@@ -1,5 +1,14 @@
 2023-08-10  Julian Brown  <julian@codesourcery.com>
 
+	* c-common.h (c_omp_region_type): Add C_ORT_UPDATE and C_ORT_OMP_UPDATE
+	codes.
+	* c-omp.cc (omp_basic_map_kind_name): New function.
+	(omp_instantiate_mapper): Add LOC parameter.  Add 'target update'
+	support.
+	(c_omp_instantiate_mappers): Add 'target update' support.
+
+2023-08-10  Julian Brown  <julian@codesourcery.com>
+
 	* c-common.h (c_omp_region_type): Add C_ORT_DECLARE_MAPPER and
 	C_ORT_OMP_DECLARE_MAPPER codes.
 
diff --git a/gcc/c-family/c-common.h b/gcc/c-family/c-common.h
index 079d1eaafaa..0998d875e3f 100644
--- a/gcc/c-family/c-common.h
+++ b/gcc/c-family/c-common.h
@@ -1271,10 +1271,12 @@ enum c_omp_region_type
   C_ORT_DECLARE_SIMD		= 1 << 2,
   C_ORT_TARGET			= 1 << 3,
   C_ORT_EXIT_DATA		= 1 << 4,
+  C_ORT_UPDATE			= 1 << 5,
   C_ORT_DECLARE_MAPPER		= 1 << 6,
   C_ORT_OMP_DECLARE_SIMD	= C_ORT_OMP | C_ORT_DECLARE_SIMD,
   C_ORT_OMP_TARGET		= C_ORT_OMP | C_ORT_TARGET,
   C_ORT_OMP_EXIT_DATA		= C_ORT_OMP | C_ORT_EXIT_DATA,
+  C_ORT_OMP_UPDATE		= C_ORT_OMP | C_ORT_UPDATE,
   C_ORT_OMP_DECLARE_MAPPER	= C_ORT_OMP | C_ORT_DECLARE_MAPPER,
   C_ORT_ACC_TARGET		= C_ORT_ACC | C_ORT_TARGET
 };
diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc
index 2ddb30b54f8..9ff09b59bc6 100644
--- a/gcc/c-family/c-omp.cc
+++ b/gcc/c-family/c-omp.cc
@@ -5176,12 +5176,37 @@ omp_map_decayed_kind (enum gomp_map_kind mapper_kind,
   return omp_join_map_kind (decay_to, force_p, always_p, present_p);
 }
 
+/* Return a name to use for a "basic" map kind, e.g. as output from
+   omp_split_map_kind above.  */
+
+static const char *
+omp_basic_map_kind_name (enum gomp_map_kind kind)
+{
+  switch (kind)
+    {
+    case GOMP_MAP_ALLOC:
+      return "alloc";
+    case GOMP_MAP_TO:
+      return "to";
+    case GOMP_MAP_FROM:
+      return "from";
+    case GOMP_MAP_TOFROM:
+      return "tofrom";
+    case GOMP_MAP_RELEASE:
+      return "release";
+    case GOMP_MAP_DELETE:
+      return "delete";
+    default:
+      gcc_unreachable ();
+    }
+}
+
 /* Instantiate a mapper MAPPER for expression EXPR, adding new clauses to
    OUTLIST.  OUTER_KIND is the mapping kind to use if not already specified in
    the mapper declaration.  */
 
 static tree *
-omp_instantiate_mapper (tree *outlist, tree mapper, tree expr,
+omp_instantiate_mapper (location_t loc, tree *outlist, tree mapper, tree expr,
 			enum gomp_map_kind outer_kind,
 			enum c_omp_region_type ort)
 {
@@ -5214,7 +5239,6 @@ omp_instantiate_mapper (tree *outlist, tree mapper, tree expr,
 
       if (TREE_CODE (t) == OMP_ARRAY_SECTION)
 	{
-	  location_t loc = OMP_CLAUSE_LOCATION (c);
 	  tree t2 = lang_hooks.decls.omp_map_array_section (loc, t);
 
 	  if (t2 == t)
@@ -5240,9 +5264,13 @@ omp_instantiate_mapper (tree *outlist, tree mapper, tree expr,
 
       walk_tree (&unshared, remap_mapper_decl_1, &map_info, NULL);
 
+      OMP_CLAUSE_LOCATION (unshared) = loc;
+
       enum gomp_map_kind decayed_kind
 	= omp_map_decayed_kind (clause_kind, outer_kind,
-				(ort & C_ORT_EXIT_DATA) != 0);
+				(ort & C_ORT_EXIT_DATA) != 0
+				|| (outer_kind == GOMP_MAP_FROM
+				    && (ort & C_ORT_UPDATE) != 0));
       OMP_CLAUSE_SET_MAP_KIND (unshared, decayed_kind);
 
       type = TYPE_MAIN_VARIANT (type);
@@ -5260,7 +5288,7 @@ omp_instantiate_mapper (tree *outlist, tree mapper, tree expr,
 	    = lang_hooks.decls.omp_extract_mapper_directive (mapper_fn);
 	  if (nested_mapper != mapper)
 	    {
-	      outlist = omp_instantiate_mapper (outlist, nested_mapper,
+	      outlist = omp_instantiate_mapper (loc, outlist, nested_mapper,
 						t, outer_kind, ort);
 	      continue;
 	    }
@@ -5271,8 +5299,51 @@ omp_instantiate_mapper (tree *outlist, tree mapper, tree expr,
 	  continue;
 	}
 
-      *outlist = unshared;
-      outlist = &OMP_CLAUSE_CHAIN (unshared);
+      if (ort & C_ORT_UPDATE)
+	{
+	  bool force_p, always_p, present_p;
+	  decayed_kind
+	    = omp_split_map_kind (decayed_kind, &force_p, &always_p,
+				  &present_p);
+	  /* We don't expect to see these flags here.  */
+	  gcc_assert (!force_p && !always_p);
+	  /* For a "target update" operation, we want to turn the map node
+	     expanded from the mapper back into a OMP_CLAUSE_TO or
+	     OMP_CLAUSE_FROM node.  If we can do neither, emit a warning and
+	     drop the clause.  */
+	  switch (decayed_kind)
+	    {
+	    case GOMP_MAP_TO:
+	    case GOMP_MAP_FROM:
+	      {
+		tree xfer
+		  = build_omp_clause (loc, (decayed_kind == GOMP_MAP_TO
+					    ? OMP_CLAUSE_TO : OMP_CLAUSE_FROM));
+		OMP_CLAUSE_DECL (xfer) = OMP_CLAUSE_DECL (unshared);
+		OMP_CLAUSE_SIZE (xfer) = OMP_CLAUSE_SIZE (unshared);
+		/* For FROM/TO clauses, "present" is represented by a flag.
+		   Set it for the expanded clause here.  */
+		if (present_p)
+		  OMP_CLAUSE_MOTION_PRESENT (xfer) = 1;
+		*outlist = xfer;
+		outlist = &OMP_CLAUSE_CHAIN (xfer);
+	      }
+	      break;
+	    default:
+	      clause_kind
+		= omp_split_map_kind (clause_kind, &force_p, &always_p,
+				      &present_p);
+	      warning_at (loc, 0, "dropping %qs clause during mapper expansion "
+			  "in %<#pragma omp target update%>",
+			  omp_basic_map_kind_name (clause_kind));
+	      inform (OMP_CLAUSE_LOCATION (c), "for map clause here");
+	    }
+	}
+      else
+	{
+	  *outlist = unshared;
+	  outlist = &OMP_CLAUSE_CHAIN (unshared);
+	}
     }
 
   return outlist;
@@ -5290,17 +5361,25 @@ c_omp_instantiate_mappers (tree clauses, enum c_omp_region_type ort)
   for (pc = &clauses, c = clauses; c; c = *pc)
     {
       bool using_mapper = false;
+      bool update_p = false, update_present_p = false;
 
       switch (OMP_CLAUSE_CODE (c))
 	{
+	case OMP_CLAUSE_TO:
+	case OMP_CLAUSE_FROM:
+	  update_p = true;
+	  if (OMP_CLAUSE_MOTION_PRESENT (c))
+	    update_present_p = true;
+	  /* Fallthrough.  */
 	case OMP_CLAUSE_MAP:
 	  {
 	    tree t = OMP_CLAUSE_DECL (c);
 	    tree type = NULL_TREE;
 	    bool nonunit_array_with_mapper = false;
 
-	    if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PUSH_MAPPER_NAME
-		|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POP_MAPPER_NAME)
+	    if (!update_p
+		&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PUSH_MAPPER_NAME
+		    || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POP_MAPPER_NAME))
 	      {
 		if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PUSH_MAPPER_NAME)
 		  mapper_name = OMP_CLAUSE_DECL (c);
@@ -5337,9 +5416,22 @@ c_omp_instantiate_mappers (tree clauses, enum c_omp_region_type ort)
 		continue;
 	      }
 
-	    enum gomp_map_kind kind = OMP_CLAUSE_MAP_KIND (c);
-	    if (kind == GOMP_MAP_UNSET)
-	      kind = GOMP_MAP_TOFROM;
+	    enum gomp_map_kind kind;
+	    if (update_p)
+	      {
+		if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO)
+		  kind = update_present_p ? GOMP_MAP_PRESENT_TO
+					  : GOMP_MAP_TO;
+		else
+		  kind = update_present_p ? GOMP_MAP_PRESENT_FROM
+					  : GOMP_MAP_FROM;
+	      }
+	    else
+	      {
+		kind = OMP_CLAUSE_MAP_KIND (c);
+		if (kind == GOMP_MAP_UNSET)
+		  kind = GOMP_MAP_TOFROM;
+	      }
 
 	    type = TYPE_MAIN_VARIANT (type);
 
@@ -5356,7 +5448,8 @@ c_omp_instantiate_mappers (tree clauses, enum c_omp_region_type ort)
 	      {
 		tree mapper
 		  = lang_hooks.decls.omp_extract_mapper_directive (mapper_fn);
-		pc = omp_instantiate_mapper (pc, mapper, t, kind, ort);
+		pc = omp_instantiate_mapper (OMP_CLAUSE_LOCATION (c),
+					     pc, mapper, t, kind, ort);
 		using_mapper = true;
 	      }
 	    else if (mapper_name)