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.
	(gfc_omp_instantiate_mapper): Add code argument to get proper
	location for diagnostic.
	(gfc_omp_instantiate_mappers): Adjust for above change.
	* 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.

Co-Authored-By: Andrew Stubbs  <ams@baylibre.com>
Co-Authored-By: Kwok Cheung Yeung <kcyeung@baylibre.com>
Co-Authored-By: Sandra Loosemore <sloosemore@baylibre.com>

1 files changed, 140 insertions, 12 deletions

diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 708f7d8..44b00ae 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -16572,7 +16572,9 @@ c_parser_omp_variable_list (c_parser *parser,
 	  bool save_c_omp_array_shaping_op_p = c_omp_array_shaping_op_p;
 	  c_omp_array_section_p = true;
 	  c_omp_array_shaping_op_p
-	    = (kind == OMP_CLAUSE_TO || kind == OMP_CLAUSE_FROM);
+	    = (kind == OMP_CLAUSE_TO
+	       || kind == OMP_CLAUSE_FROM
+	       || ort == C_ORT_OMP_DECLARE_MAPPER);
 	  c_expr expr = c_parser_expr_no_commas (parser, NULL);
 	  if (expr.value != error_mark_node)
 	    mark_exp_read (expr.value);
@@ -20586,7 +20588,9 @@ c_parser_omp_clause_map (c_parser *parser, tree list, enum gomp_map_kind kind)
     }
 
   nl = c_parser_omp_variable_list (parser, clause_loc, OMP_CLAUSE_MAP, list,
-				   C_ORT_OMP, true);
+				   (kind == GOMP_MAP_UNSET
+				    ? C_ORT_OMP_DECLARE_MAPPER
+				    : C_ORT_OMP), true);
 
   tree last_new = NULL_TREE;
 
@@ -20863,25 +20867,147 @@ c_parser_omp_clause_from_to (c_parser *parser, enum omp_clause_code kind,
   if (!parens.require_open (parser))
     return list;
 
-  bool present = false;
-  c_token *token = c_parser_peek_token (parser);
+  int pos = 1, colon_pos = 0;
 
-  if (token->type == CPP_NAME
-      && strcmp (IDENTIFIER_POINTER (token->value), "present") == 0
-      && c_parser_peek_2nd_token (parser)->type == CPP_COLON)
+  while (c_parser_peek_nth_token_raw (parser, pos)->type == CPP_NAME)
     {
-      present = true;
-      c_parser_consume_token (parser);
-      c_parser_consume_token (parser);
+      if (c_parser_peek_nth_token_raw (parser, pos + 1)->type == CPP_COMMA)
+	pos += 2;
+      else if (c_parser_peek_nth_token_raw (parser, pos + 1)->type
+	       == CPP_OPEN_PAREN)
+	{
+	  unsigned int npos = pos + 2;
+	  if (c_parser_check_balanced_raw_token_sequence (parser, &npos)
+	     && (c_parser_peek_nth_token_raw (parser, npos)->type
+		 == CPP_CLOSE_PAREN))
+	    pos = npos + 1;
+	}
+      else
+	pos++;
+      if (c_parser_peek_nth_token_raw (parser, pos)->type == CPP_COLON)
+	{
+	  colon_pos = pos;
+	  break;
+	}
     }
 
+  int present_modifier = false;
+  int mapper_modifier = false;
+  tree mapper_name = NULL_TREE;
+
+  for (int pos = 1; pos < colon_pos; ++pos)
+    {
+      c_token *tok = c_parser_peek_token (parser);
+      if (tok->type == CPP_COMMA)
+	{
+	  c_parser_consume_token (parser);
+	  continue;
+	}
+      const char *p = IDENTIFIER_POINTER (tok->value);
+      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++;
+	  c_parser_consume_token (parser);
+	}
+      else if (strcmp ("mapper", p) == 0)
+	{
+	  c_parser_consume_token (parser);
+
+	  matching_parens mparens;
+	  if (mparens.require_open (parser))
+	    {
+	      if (mapper_modifier)
+		{
+		  c_parser_error (parser, "too many %<mapper%> modifiers");
+		  /* Assume it's a well-formed mapper modifier, even if it
+		     seems to be in the wrong place.  */
+		  c_parser_consume_token (parser);
+		  mparens.require_close (parser);
+		  parens.skip_until_found_close (parser);
+		  return list;
+		}
+
+	      tok = c_parser_peek_token (parser);
+
+	      switch (tok->type)
+		{
+		case CPP_NAME:
+		  {
+		    mapper_name = tok->value;
+		    c_parser_consume_token (parser);
+		  }
+		  break;
+
+		case CPP_KEYWORD:
+		  if (tok->keyword == RID_DEFAULT)
+		    {
+		      c_parser_consume_token (parser);
+		      break;
+		    }
+		  /* Fallthrough.  */
+
+		default:
+		  error_at (tok->location,
+			    "expected identifier or %<default%>");
+		  return list;
+		}
+
+	      if (!mparens.require_close (parser))
+		{
+		  parens.skip_until_found_close (parser);
+		  return list;
+		}
+
+	      mapper_modifier++;
+	      pos += 3;
+	    }
+	}
+      else
+	{
+	  c_parser_error (parser, "%<to%> or %<from%> clause with modifier "
+			  "other than %<present%> or %<mapper%>");
+	  parens.skip_until_found_close (parser);
+	  return list;
+	}
+    }
+
+  if (colon_pos)
+    c_parser_require (parser, CPP_COLON, "expected %<:%>");
+
   tree nl = c_parser_omp_variable_list (parser, loc, kind, list, C_ORT_OMP, true);
   parens.skip_until_found_close (parser);
 
-  if (present)
+  if (present_modifier)
     for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
       OMP_CLAUSE_MOTION_PRESENT (c) = 1;
 
+  if (mapper_name)
+    {
+      tree last_new = NULL_TREE;
+      for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
+	last_new = c;
+
+      tree name = build_omp_clause (input_location, OMP_CLAUSE_MAP);
+      OMP_CLAUSE_SET_MAP_KIND (name, GOMP_MAP_PUSH_MAPPER_NAME);
+      OMP_CLAUSE_DECL (name) = mapper_name;
+      OMP_CLAUSE_CHAIN (name) = nl;
+      nl = name;
+
+      gcc_assert (last_new);
+
+      name = build_omp_clause (input_location, OMP_CLAUSE_MAP);
+      OMP_CLAUSE_SET_MAP_KIND (name, GOMP_MAP_POP_MAPPER_NAME);
+      OMP_CLAUSE_DECL (name) = null_pointer_node;
+      OMP_CLAUSE_CHAIN (name) = OMP_CLAUSE_CHAIN (last_new);
+      OMP_CLAUSE_CHAIN (last_new) = name;
+    }
+
   return nl;
 }
 
@@ -26536,7 +26662,9 @@ c_parser_omp_target_update (location_t loc, c_parser *parser,
 
   tree clauses
     = c_parser_omp_all_clauses (parser, OMP_TARGET_UPDATE_CLAUSE_MASK,
-				"#pragma omp target update");
+				"#pragma omp target update", false);
+  clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP_UPDATE);
+  clauses = c_finish_omp_clauses (clauses, C_ORT_OMP_UPDATE);
   bool to_clause = false, from_clause = false;
   for (tree c = clauses;
        c && !to_clause && !from_clause;