[PATCH 4/8] OpenMP: Support OpenMP 5.0 "declare mapper" directives for C

Julian Brown julian@codesourcery.com
Tue Sep 5 19:28:24 GMT 2023


This patch adds support for "declare mapper" directives (and the "mapper"
modifier on "map" clauses) for C.  It was previously posted for mainline
here:

  https://gcc.gnu.org/pipermail/gcc-patches/2022-December/609041.html

and for the og13 branch here:

  https://gcc.gnu.org/pipermail/gcc-patches/2023-June/623356.html

This version supports mappers on "target data", "target enter data" and
"target exit data" directives as well as on "target" directives.

2023-09-05  Julian Brown  <julian@codesourcery.com>

gcc/c/
	* c-decl.cc (c_omp_mapper_id, c_omp_mapper_decl, c_omp_mapper_lookup,
	c_omp_extract_mapper_directive, c_omp_map_array_section,
	c_omp_scan_mapper_bindings_r, c_omp_scan_mapper_bindings): New
	functions.
	* c-objc-common.h (LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES,
	LANG_HOOKS_OMP_MAPPER_LOOKUP, LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE,
	LANG_HOOKS_OMP_MAP_ARRAY_SECTION): Define langhooks for C.
	* c-parser.cc (c_parser_omp_clause_map): Add KIND parameter.  Handle
	mapper modifier.
	(c_parser_omp_all_clauses): Update call to c_parser_omp_clause_map with
	new kind argument.
	(c_parser_omp_target_data, c_parser_omp_target_enter_data,
	c_parser_omp_target_exit_data): Instantiate explicit mappers.
	(c_parser_omp_target): Instantiate explicit mappers and record bindings
	for implicit mappers.
	(c_parser_omp_declare_mapper): Parse "declare mapper" directives.
	(c_parser_omp_declare): Support "declare mapper".
	* c-tree.h (c_omp_finish_mapper_clauses, c_omp_mapper_lookup,
	c_omp_extract_mapper_directive, c_omp_map_array_section,
	c_omp_mapper_id, c_omp_mapper_decl, c_omp_scan_mapper_bindings): Add
	prototypes.
	* c-typeck.cc (c_finish_omp_clauses): Handle GOMP_MAP_PUSH_MAPPER_NAME
	and GOMP_MAP_POP_MAPPER_NAME.
	(c_omp_finish_mapper_clauses): New function (langhook).

gcc/testsuite/
	* c-c++-common/gomp/declare-mapper-3.c: Enable for C.
	* c-c++-common/gomp/declare-mapper-4.c: Likewise.
	* c-c++-common/gomp/declare-mapper-5.c: Likewise.
	* c-c++-common/gomp/declare-mapper-6.c: Likewise.
	* c-c++-common/gomp/declare-mapper-7.c: Likewise.
	* c-c++-common/gomp/declare-mapper-8.c: Likewise.
	* c-c++-common/gomp/declare-mapper-9.c: Likewise.
	* c-c++-common/gomp/declare-mapper-12.c: Likewise.
	* c-c++-common/gomp/declare-mapper-15.c: Likewise.
	* c-c++-common/gomp/declare-mapper-16.c: Likewise.
	* gcc.dg/gomp/declare-mapper-10.c: New test.
	* gcc.dg/gomp/declare-mapper-11.c: New test.

libgomp/
	* testsuite/libgomp.c-c++-common/declare-mapper-9.c: Enable for C.
	* testsuite/libgomp.c-c++-common/declare-mapper-10.c: Likewise.
	* testsuite/libgomp.c-c++-common/declare-mapper-11.c: Likewise.
	* testsuite/libgomp.c-c++-common/declare-mapper-12.c: Likewise.
	* testsuite/libgomp.c-c++-common/declare-mapper-13.c: Likewise.
	* testsuite/libgomp.c-c++-common/declare-mapper-14.c: Likewise.
---
 gcc/c/c-decl.cc                               | 169 ++++++++++
 gcc/c/c-objc-common.h                         |  12 +
 gcc/c/c-parser.cc                             | 291 ++++++++++++++++--
 gcc/c/c-tree.h                                |   7 +
 gcc/c/c-typeck.cc                             |  15 +
 .../c-c++-common/gomp/declare-mapper-12.c     |   2 +-
 .../c-c++-common/gomp/declare-mapper-15.c     |   2 +-
 .../c-c++-common/gomp/declare-mapper-16.c     |   2 +-
 .../c-c++-common/gomp/declare-mapper-3.c      |   2 +-
 .../c-c++-common/gomp/declare-mapper-4.c      |   2 +-
 .../c-c++-common/gomp/declare-mapper-5.c      |   2 +-
 .../c-c++-common/gomp/declare-mapper-6.c      |   2 +-
 .../c-c++-common/gomp/declare-mapper-7.c      |   2 +-
 .../c-c++-common/gomp/declare-mapper-8.c      |   2 +-
 .../c-c++-common/gomp/declare-mapper-9.c      |   2 +-
 gcc/testsuite/gcc.dg/gomp/declare-mapper-10.c |  61 ++++
 gcc/testsuite/gcc.dg/gomp/declare-mapper-11.c |  33 ++
 .../libgomp.c-c++-common/declare-mapper-10.c  |   2 +-
 .../libgomp.c-c++-common/declare-mapper-11.c  |   2 +-
 .../libgomp.c-c++-common/declare-mapper-12.c  |   2 +-
 .../libgomp.c-c++-common/declare-mapper-13.c  |   2 +-
 .../libgomp.c-c++-common/declare-mapper-14.c  |   2 +-
 .../libgomp.c-c++-common/declare-mapper-9.c   |   2 +-
 23 files changed, 584 insertions(+), 36 deletions(-)
 create mode 100644 gcc/testsuite/gcc.dg/gomp/declare-mapper-10.c
 create mode 100644 gcc/testsuite/gcc.dg/gomp/declare-mapper-11.c

diff --git a/gcc/c/c-decl.cc b/gcc/c/c-decl.cc
index 1f9eb44dbaa2..ac71092fcad6 100644
--- a/gcc/c/c-decl.cc
+++ b/gcc/c/c-decl.cc
@@ -13144,6 +13144,175 @@ c_check_omp_declare_reduction_r (tree *tp, int *, void *data)
   return NULL_TREE;
 }
 
+/* Return identifier to look up for omp declare reduction.  */
+
+tree
+c_omp_mapper_id (tree mapper_id)
+{
+  const char *p = NULL;
+
+  const char prefix[] = "omp declare mapper ";
+
+  if (mapper_id == NULL_TREE)
+    p = "<default>";
+  else if (TREE_CODE (mapper_id) == IDENTIFIER_NODE)
+    p = IDENTIFIER_POINTER (mapper_id);
+  else
+    return error_mark_node;
+
+  size_t lenp = sizeof (prefix);
+  size_t len = strlen (p);
+  char *name = XALLOCAVEC (char, lenp + len);
+  memcpy (name, prefix, lenp - 1);
+  memcpy (name + lenp - 1, p, len + 1);
+  return get_identifier (name);
+}
+
+/* Lookup MAPPER_ID in the current scope, or create an artificial
+   VAR_DECL, bind it into the current scope and return it.  */
+
+tree
+c_omp_mapper_decl (tree mapper_id)
+{
+  struct c_binding *b = I_SYMBOL_BINDING (mapper_id);
+  if (b != NULL && B_IN_CURRENT_SCOPE (b))
+    return b->decl;
+
+  tree decl = build_decl (BUILTINS_LOCATION, VAR_DECL,
+			  mapper_id, integer_type_node);
+  DECL_ARTIFICIAL (decl) = 1;
+  DECL_EXTERNAL (decl) = 1;
+  TREE_STATIC (decl) = 1;
+  TREE_PUBLIC (decl) = 0;
+  bind (mapper_id, decl, current_scope, true, false, BUILTINS_LOCATION);
+  return decl;
+}
+
+/* Lookup MAPPER_ID in the first scope where it has entry for TYPE.  */
+
+tree
+c_omp_mapper_lookup (tree mapper_id, tree type)
+{
+  if (TREE_CODE (type) != RECORD_TYPE
+      && TREE_CODE (type) != UNION_TYPE)
+    return NULL_TREE;
+
+  mapper_id = c_omp_mapper_id (mapper_id);
+
+  struct c_binding *b = I_SYMBOL_BINDING (mapper_id);
+  while (b)
+    {
+      tree t;
+      for (t = DECL_INITIAL (b->decl); t; t = TREE_CHAIN (t))
+	if (comptypes (TREE_PURPOSE (t), type))
+	  return TREE_VALUE (t);
+      b = b->shadowed;
+    }
+  return NULL_TREE;
+}
+
+/* For C, we record a pointer to the mapper itself without wrapping it in an
+   artificial function or similar.  So, just return it.  */
+
+tree
+c_omp_extract_mapper_directive (tree mapper)
+{
+  return mapper;
+}
+
+/* For now we can handle singleton OMP_ARRAY_SECTIONs with custom mappers, but
+   nothing more complicated.  */
+
+tree
+c_omp_map_array_section (location_t loc, tree t)
+{
+  tree low = TREE_OPERAND (t, 1);
+  tree len = TREE_OPERAND (t, 2);
+
+  if (len && integer_onep (len))
+    {
+      t = TREE_OPERAND (t, 0);
+
+      if (!low)
+	low = integer_zero_node;
+
+      t = build_array_ref (loc, t, low);
+    }
+
+  return t;
+}
+
+/* Helper function for below function.  */
+
+static tree
+c_omp_scan_mapper_bindings_r (tree *tp, int *walk_subtrees, void *ptr)
+{
+  tree t = *tp;
+  omp_mapper_list<tree> *mlist = (omp_mapper_list<tree> *) ptr;
+  tree aggr_type = NULL_TREE;
+
+  if (TREE_CODE (t) == SIZEOF_EXPR
+      || TREE_CODE (t) == ALIGNOF_EXPR)
+    {
+      *walk_subtrees = 0;
+      return NULL_TREE;
+    }
+
+  if (TREE_CODE (t) == OMP_CLAUSE)
+    return NULL_TREE;
+
+  if (TREE_CODE (t) == COMPONENT_REF
+      && AGGREGATE_TYPE_P (TREE_TYPE (TREE_OPERAND (t, 0))))
+    aggr_type = TREE_TYPE (TREE_OPERAND (t, 0));
+  else if ((TREE_CODE (t) == VAR_DECL
+	    || TREE_CODE (t) == PARM_DECL
+	    || TREE_CODE (t) == RESULT_DECL)
+	   && AGGREGATE_TYPE_P (TREE_TYPE (t)))
+    aggr_type = TREE_TYPE (t);
+
+  if (aggr_type)
+    {
+      tree mapper_fn = c_omp_mapper_lookup (NULL_TREE, aggr_type);
+      if (mapper_fn)
+	mlist->add_mapper (NULL_TREE, aggr_type, mapper_fn);
+    }
+
+  return NULL_TREE;
+}
+
+/* Scan an offload region's body, and record uses of struct- or union-typed
+   variables.  Add _mapper_binding_ fake clauses to *CLAUSES_PTR.  */
+
+void
+c_omp_scan_mapper_bindings (location_t loc, tree *clauses_ptr, tree body)
+{
+  hash_set<omp_name_type<tree>> seen_types;
+  auto_vec<tree> mappers;
+  omp_mapper_list<tree> mlist (&seen_types, &mappers);
+
+  walk_tree_without_duplicates (&body, c_omp_scan_mapper_bindings_r, &mlist);
+
+  unsigned int i;
+  tree mapper;
+  FOR_EACH_VEC_ELT (mappers, i, mapper)
+    c_omp_find_nested_mappers (&mlist, mapper);
+
+  FOR_EACH_VEC_ELT (mappers, i, mapper)
+    {
+      if (mapper == error_mark_node)
+	continue;
+      tree mapper_name = OMP_DECLARE_MAPPER_ID (mapper);
+      tree decl = OMP_DECLARE_MAPPER_DECL (mapper);
+
+      tree c = build_omp_clause (loc, OMP_CLAUSE__MAPPER_BINDING_);
+      OMP_CLAUSE__MAPPER_BINDING__ID (c) = mapper_name;
+      OMP_CLAUSE__MAPPER_BINDING__DECL (c) = decl;
+      OMP_CLAUSE__MAPPER_BINDING__MAPPER (c) = mapper;
+
+      OMP_CLAUSE_CHAIN (c) = *clauses_ptr;
+      *clauses_ptr = c;
+    }
+}
 
 bool
 c_check_in_current_scope (tree decl)
diff --git a/gcc/c/c-objc-common.h b/gcc/c/c-objc-common.h
index d31dacb9dd47..423b0d3f434c 100644
--- a/gcc/c/c-objc-common.h
+++ b/gcc/c/c-objc-common.h
@@ -122,6 +122,18 @@ along with GCC; see the file COPYING3.  If not see
 #undef LANG_HOOKS_OMP_CLAUSE_ASSIGN_OP
 #define LANG_HOOKS_OMP_CLAUSE_ASSIGN_OP c_omp_clause_copy_ctor
 
+#undef LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES
+#define LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES c_omp_finish_mapper_clauses
+
+#undef LANG_HOOKS_OMP_MAPPER_LOOKUP
+#define LANG_HOOKS_OMP_MAPPER_LOOKUP c_omp_mapper_lookup
+
+#undef LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE
+#define LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE c_omp_extract_mapper_directive
+
+#undef LANG_HOOKS_OMP_MAP_ARRAY_SECTION
+#define LANG_HOOKS_OMP_MAP_ARRAY_SECTION c_omp_map_array_section
+
 #undef LANG_HOOKS_TREE_INLINING_VAR_MOD_TYPE_P
 #define LANG_HOOKS_TREE_INLINING_VAR_MOD_TYPE_P c_var_mod_p
 #endif /* GCC_C_OBJC_COMMON */
diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 36215096af4c..35b9ba6f2f0c 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -17321,10 +17321,9 @@ c_parser_omp_clause_doacross (c_parser *parser, tree list)
      always | close */
 
 static tree
-c_parser_omp_clause_map (c_parser *parser, tree list)
+c_parser_omp_clause_map (c_parser *parser, tree list, enum gomp_map_kind kind)
 {
   location_t clause_loc = c_parser_peek_token (parser)->location;
-  enum gomp_map_kind kind = GOMP_MAP_TOFROM;
   tree nl, c;
 
   matching_parens parens;
@@ -17343,12 +17342,26 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
 
       if (c_parser_peek_nth_token_raw (parser, pos + 1)->type == CPP_COMMA)
 	pos++;
+      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)
+	      && (c_parser_peek_nth_token_raw (parser, npos + 1)->type
+		  == CPP_COMMA))
+	    pos = npos + 1;
+	}
+
       pos++;
     }
 
   int always_modifier = 0;
   int close_modifier = 0;
   int present_modifier = 0;
+  int mapper_modifier = 0;
+  tree mapper_name = NULL_TREE;
   for (int pos = 1; pos < map_kind_pos; ++pos)
     {
       c_token *tok = c_parser_peek_token (parser);
@@ -17369,6 +17382,7 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
 	      return list;
 	    }
 	  always_modifier++;
+	  c_parser_consume_token (parser);
 	}
       else if (strcmp ("close", p) == 0)
 	{
@@ -17379,6 +17393,60 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
 	      return list;
 	    }
 	  close_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 if (strcmp ("present", p) == 0)
 	{
@@ -17389,16 +17457,16 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
 	      return list;
 	    }
 	  present_modifier++;
+	  c_parser_consume_token (parser);
 	}
       else
 	{
 	  c_parser_error (parser, "%<map%> clause with map-type modifier other "
-				  "than %<always%>, %<close%> or %<present%>");
+				  "than %<always%>, %<close%>, %<mapper%> or "
+				  "%<present%>");
 	  parens.skip_until_found_close (parser);
 	  return list;
 	}
-
-	c_parser_consume_token (parser);
     }
 
   if (c_parser_next_token_is (parser, CPP_NAME)
@@ -17442,8 +17510,30 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
   nl = c_parser_omp_variable_list (parser, clause_loc, OMP_CLAUSE_MAP, list,
 				   true);
 
+  tree last_new = NULL_TREE;
+
   for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
-    OMP_CLAUSE_SET_MAP_KIND (c, kind);
+    {
+      OMP_CLAUSE_SET_MAP_KIND (c, kind);
+      last_new = c;
+    }
+
+  if (mapper_name)
+    {
+      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;
+    }
 
   parens.skip_until_found_close (parser);
   return nl;
@@ -18269,7 +18359,7 @@ c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask,
 	  c_name = "doacross";
 	  break;
 	case PRAGMA_OMP_CLAUSE_MAP:
-	  clauses = c_parser_omp_clause_map (parser, clauses);
+	  clauses = c_parser_omp_clause_map (parser, clauses, GOMP_MAP_TOFROM);
 	  c_name = "map";
 	  break;
 	case PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR:
@@ -22051,7 +22141,9 @@ c_parser_omp_target_data (location_t loc, c_parser *parser, bool *if_p)
 
   tree clauses
     = c_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK,
-				"#pragma omp target data");
+				"#pragma omp target data", false);
+  clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP);
+  clauses = c_finish_omp_clauses (clauses, C_ORT_OMP);
   c_omp_adjust_map_clauses (clauses, false);
   int map_seen = 0;
   for (tree *pc = &clauses; *pc;)
@@ -22208,7 +22300,9 @@ c_parser_omp_target_enter_data (location_t loc, c_parser *parser,
 
   tree clauses
     = c_parser_omp_all_clauses (parser, OMP_TARGET_ENTER_DATA_CLAUSE_MASK,
-				"#pragma omp target enter data");
+				"#pragma omp target enter data", false);
+  clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP);
+  clauses = c_finish_omp_clauses (clauses, C_ORT_OMP);
   c_omp_adjust_map_clauses (clauses, false);
   int map_seen = 0;
   for (tree *pc = &clauses; *pc;)
@@ -22318,7 +22412,9 @@ c_parser_omp_target_exit_data (location_t loc, c_parser *parser,
 
   tree clauses
     = c_parser_omp_all_clauses (parser, OMP_TARGET_EXIT_DATA_CLAUSE_MASK,
-				"#pragma omp target exit data");
+				"#pragma omp target exit data", false);
+  clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP_EXIT_DATA);
+  clauses = c_finish_omp_clauses (clauses, C_ORT_OMP_EXIT_DATA);
   c_omp_adjust_map_clauses (clauses, false);
   int map_seen = 0;
   for (tree *pc = &clauses; *pc;)
@@ -22408,7 +22504,7 @@ c_parser_omp_target (c_parser *parser, enum pragma_context context, bool *if_p)
 {
   location_t loc = c_parser_peek_token (parser)->location;
   c_parser_consume_pragma (parser);
-  tree *pc = NULL, stmt, block;
+  tree *pc = NULL, stmt, block, body, clauses;
 
   if (context != pragma_stmt && context != pragma_compound)
     {
@@ -22563,10 +22659,9 @@ c_parser_omp_target (c_parser *parser, enum pragma_context context, bool *if_p)
   stmt = make_node (OMP_TARGET);
   TREE_TYPE (stmt) = void_type_node;
 
-  OMP_TARGET_CLAUSES (stmt)
-    = c_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK,
-				"#pragma omp target", false);
-  for (tree c = OMP_TARGET_CLAUSES (stmt); c; c = OMP_CLAUSE_CHAIN (c))
+  clauses = c_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK,
+				      "#pragma omp target", false);
+  for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
     if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IN_REDUCTION)
       {
 	tree nc = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
@@ -22575,14 +22670,19 @@ c_parser_omp_target (c_parser *parser, enum pragma_context context, bool *if_p)
 	OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (c);
 	OMP_CLAUSE_CHAIN (c) = nc;
       }
-  OMP_TARGET_CLAUSES (stmt)
-    = c_finish_omp_clauses (OMP_TARGET_CLAUSES (stmt), C_ORT_OMP_TARGET);
-  c_omp_adjust_map_clauses (OMP_TARGET_CLAUSES (stmt), true);
+  clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP_TARGET);
+  clauses  = c_finish_omp_clauses (clauses, C_ORT_OMP_TARGET);
+  c_omp_adjust_map_clauses (clauses, true);
 
-  pc = &OMP_TARGET_CLAUSES (stmt);
   keep_next_level ();
   block = c_begin_compound_stmt (true);
-  add_stmt (c_parser_omp_structured_block (parser, if_p));
+  body = c_parser_omp_structured_block (parser, if_p);
+
+  c_omp_scan_mapper_bindings (loc, &clauses, body);
+
+  add_stmt (body);
+  OMP_TARGET_CLAUSES (stmt) = clauses;
+  pc = &OMP_TARGET_CLAUSES (stmt);
   OMP_TARGET_BODY (stmt) = c_end_compound_stmt (loc, block, true);
 
   SET_EXPR_LOCATION (stmt, loc);
@@ -23891,6 +23991,151 @@ c_parser_omp_declare_reduction (c_parser *parser, enum pragma_context context)
 }
 
 
+/* OpenMP 5.0
+   #pragma omp declare mapper ([mapper-identifier :] type var) \
+			      [clause [ [,] clause ] ... ] new-line  */
+
+static void
+c_parser_omp_declare_mapper (c_parser *parser, enum pragma_context context)
+{
+  tree type, mapper_name = NULL_TREE, var = NULL_TREE, stmt, stmtlist;
+  tree maplist = NULL_TREE, mapper_id, mapper_decl, t;
+  c_token *token;
+
+  if (context == pragma_struct || context == pragma_param)
+    {
+      error ("%<#pragma omp declare reduction%> not at file or block scope");
+      goto fail;
+    }
+
+  if (!c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>"))
+    goto fail;
+
+  token = c_parser_peek_token (parser);
+
+  if (c_parser_peek_2nd_token (parser)->type == CPP_COLON)
+    {
+      switch (token->type)
+	{
+	case CPP_NAME:
+	  mapper_name = token->value;
+	  c_parser_consume_token (parser);
+	  break;
+	case CPP_KEYWORD:
+	  if (token->keyword == RID_DEFAULT)
+	    {
+	      mapper_name = NULL_TREE;
+	      c_parser_consume_token (parser);
+	      break;
+	    }
+	  /* Fallthrough.  */
+	default:
+	  error_at (token->location, "expected identifier or %<default%>");
+	  c_parser_skip_to_pragma_eol (parser, false);
+	  return;
+	}
+
+      if (!c_parser_require (parser, CPP_COLON, "expected %<:%>"))
+	goto fail;
+    }
+
+  mapper_id = c_omp_mapper_id (mapper_name);
+  mapper_decl = c_omp_mapper_decl (mapper_id);
+
+  {
+    location_t loc = c_parser_peek_token (parser)->location;
+    struct c_type_name *ctype = c_parser_type_name (parser);
+    type = groktypename (ctype, NULL, NULL);
+    if (type == error_mark_node)
+      goto fail;
+    if (TREE_CODE (type) != RECORD_TYPE
+	&& TREE_CODE (type) != UNION_TYPE)
+      {
+	error_at (loc, "%qT is not a struct or union type in "
+		  "%<#pragma omp declare mapper%>", type);
+	c_parser_skip_to_pragma_eol (parser, false);
+	return;
+      }
+    for (tree t = DECL_INITIAL (mapper_decl); t; t = TREE_CHAIN (t))
+      if (comptypes (TREE_PURPOSE (t), type))
+	{
+	  error_at (loc, "redeclaration of %qs %<#pragma omp declare "
+		    "mapper%> for type %qT", IDENTIFIER_POINTER (mapper_id)
+		      + sizeof ("omp declare mapper ") - 1,
+		    type);
+	  tree prevmapper = TREE_VALUE (t);
+	  /* Hmm, this location might not be very accurate.  */
+	  location_t ploc
+	    = DECL_SOURCE_LOCATION (OMP_DECLARE_MAPPER_DECL (prevmapper));
+	  error_at (ploc, "previous %<#pragma omp declare mapper%>");
+	  c_parser_skip_to_pragma_eol (parser, false);
+	  return;
+	}
+  }
+
+  token = c_parser_peek_token (parser);
+  if (token->type == CPP_NAME)
+    {
+      var = build_decl (token->location, VAR_DECL, token->value, type);
+      c_parser_consume_token (parser);
+      DECL_ARTIFICIAL (var) = 1;
+    }
+  else
+    {
+      error_at (token->location, "expected identifier");
+      goto fail;
+    }
+
+  if (!c_parser_require (parser, CPP_CLOSE_PAREN, "expected %<)%>"))
+    goto fail;
+
+  push_scope ();
+  stmtlist = push_stmt_list ();
+  pushdecl (var);
+  DECL_CONTEXT (var) = current_function_decl;
+
+  while (c_parser_next_token_is_not (parser, CPP_PRAGMA_EOL))
+    {
+      location_t here;
+      pragma_omp_clause c_kind;
+      here = c_parser_peek_token (parser)->location;
+      c_kind = c_parser_omp_clause_name (parser);
+      if (c_kind != PRAGMA_OMP_CLAUSE_MAP)
+	{
+	  error_at (here, "unexpected clause");
+	  goto fail;
+	}
+      maplist = c_parser_omp_clause_map (parser, maplist, GOMP_MAP_UNSET);
+    }
+
+  if (maplist == NULL_TREE)
+    {
+      error_at (input_location, "missing %<map%> clause");
+      goto fail;
+    }
+
+  stmt = make_node (OMP_DECLARE_MAPPER);
+  TREE_TYPE (stmt) = type;
+  OMP_DECLARE_MAPPER_ID (stmt) = mapper_name;
+  OMP_DECLARE_MAPPER_DECL (stmt) = var;
+  OMP_DECLARE_MAPPER_CLAUSES (stmt) = maplist;
+
+  add_stmt (stmt);
+
+  pop_stmt_list (stmtlist);
+  pop_scope ();
+
+  c_parser_skip_to_pragma_eol (parser);
+
+  t = tree_cons (type, stmt, DECL_INITIAL (mapper_decl));
+  DECL_INITIAL (mapper_decl) = t;
+
+  return;
+
+ fail:
+  c_parser_skip_to_pragma_eol (parser);
+}
+
 /* OpenMP 4.0
    #pragma omp declare simd declare-simd-clauses[optseq] new-line
    #pragma omp declare reduction (reduction-id : typename-list : expression) \
@@ -23920,6 +24165,12 @@ c_parser_omp_declare (c_parser *parser, enum pragma_context context)
 	  c_parser_omp_declare_reduction (parser, context);
 	  return false;
 	}
+      if (strcmp (p, "mapper") == 0)
+	{
+	  c_parser_consume_token (parser);
+	  c_parser_omp_declare_mapper (parser, context);
+	  return false;
+	}
       if (!flag_openmp)  /* flag_openmp_simd  */
 	{
 	  c_parser_skip_to_pragma_eol (parser, false);
diff --git a/gcc/c/c-tree.h b/gcc/c/c-tree.h
index 5cd507563695..d2cc975a5ba3 100644
--- a/gcc/c/c-tree.h
+++ b/gcc/c/c-tree.h
@@ -826,6 +826,10 @@ extern tree c_finish_omp_task (location_t, tree, tree);
 extern void c_finish_omp_cancel (location_t, tree);
 extern void c_finish_omp_cancellation_point (location_t, tree);
 extern tree c_finish_omp_clauses (tree, enum c_omp_region_type);
+extern tree c_omp_finish_mapper_clauses (tree);
+extern tree c_omp_mapper_lookup (tree, tree);
+extern tree c_omp_extract_mapper_directive (tree);
+extern tree c_omp_map_array_section (location_t, tree);
 extern tree c_build_va_arg (location_t, tree, location_t, tree);
 extern tree c_finish_transaction (location_t, tree, int);
 extern bool c_tree_equal (tree, tree);
@@ -877,6 +881,9 @@ extern tree c_omp_reduction_id (enum tree_code, tree);
 extern tree c_omp_reduction_decl (tree);
 extern tree c_omp_reduction_lookup (tree, tree);
 extern tree c_check_omp_declare_reduction_r (tree *, int *, void *);
+extern tree c_omp_mapper_id (tree);
+extern tree c_omp_mapper_decl (tree);
+extern void c_omp_scan_mapper_bindings (location_t, tree *, tree);
 extern bool c_check_in_current_scope (tree);
 extern void c_pushtag (location_t, tree, tree);
 extern void c_bind (location_t, tree, bool);
diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc
index b399341084d2..2696e681be4f 100644
--- a/gcc/c/c-typeck.cc
+++ b/gcc/c/c-typeck.cc
@@ -15336,6 +15336,12 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	case OMP_CLAUSE_MAP:
 	  if (OMP_CLAUSE_MAP_IMPLICIT (c) && !implicit_moved)
 	    goto move_implicit;
+	  if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PUSH_MAPPER_NAME
+	      || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POP_MAPPER_NAME)
+	    {
+	      remove = true;
+	      break;
+	    }
 	  /* FALLTHRU */
 	case OMP_CLAUSE_TO:
 	case OMP_CLAUSE_FROM:
@@ -16180,6 +16186,15 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
   return clauses;
 }
 
+/* Do processing necessary to make CLAUSES well-formed, where CLAUSES result
+   from implicit instantiation of user-defined mappers (in gimplify.cc).  */
+
+tree
+c_omp_finish_mapper_clauses (tree clauses)
+{
+  return c_finish_omp_clauses (clauses, C_ORT_OMP);
+}
+
 /* Return code to initialize DST with a copy constructor from SRC.
    C doesn't have copy constructors nor assignment operators, only for
    _Atomic vars we need to perform __atomic_load from src into a temporary
diff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-12.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-12.c
index c4d017036c5e..dffb19db03cd 100644
--- a/gcc/testsuite/c-c++-common/gomp/declare-mapper-12.c
+++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-12.c
@@ -1,4 +1,4 @@
-/* { dg-do compile { target c++ } } */
+/* { dg-do compile } */
 
 struct XYZ {
   int a;
diff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-15.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-15.c
index 0db83b6fd335..ecda2e5ebd1a 100644
--- a/gcc/testsuite/c-c++-common/gomp/declare-mapper-15.c
+++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-15.c
@@ -1,4 +1,4 @@
-/* { dg-do compile { target c++ } } */
+/* { dg-do compile } */
 /* { dg-options "-fopenmp -fdump-tree-gimple" } */
 
 typedef struct {
diff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-16.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-16.c
index f8711f60a39a..20383cc2d69f 100644
--- a/gcc/testsuite/c-c++-common/gomp/declare-mapper-16.c
+++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-16.c
@@ -1,4 +1,4 @@
-/* { dg-do compile { target c++ } } */
+/* { dg-do compile } */
 /* { dg-options "-fopenmp -fdump-tree-gimple" } */
 
 typedef struct {
diff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-3.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-3.c
index 983d979d68c5..e491bcd0ce65 100644
--- a/gcc/testsuite/c-c++-common/gomp/declare-mapper-3.c
+++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-3.c
@@ -1,4 +1,4 @@
-// { dg-do compile { target c++ } }
+// { dg-do compile }
 // { dg-additional-options "-fdump-tree-gimple" }
 
 #include <stdlib.h>
diff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-4.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-4.c
index 6d933e4bf6f4..39e3ab114199 100644
--- a/gcc/testsuite/c-c++-common/gomp/declare-mapper-4.c
+++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-4.c
@@ -1,4 +1,4 @@
-/* { dg-do compile { target c++ } } */
+/* { dg-do compile } */
 /* { dg-additional-options "-fdump-tree-original" } */
 
 /* Check mapper binding clauses.  */
diff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-5.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-5.c
index f675a8c68902..86f14e76cbf3 100644
--- a/gcc/testsuite/c-c++-common/gomp/declare-mapper-5.c
+++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-5.c
@@ -1,4 +1,4 @@
-/* { dg-do compile { target c++ } } */
+/* { dg-do compile } */
 
 typedef struct S_ {
   int *myarr;
diff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-6.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-6.c
index a2f6c08cdfdd..c13eb8b5816e 100644
--- a/gcc/testsuite/c-c++-common/gomp/declare-mapper-6.c
+++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-6.c
@@ -1,4 +1,4 @@
-/* { dg-do compile { target c++ } } */
+/* { dg-do compile } */
 
 int x = 5;
 
diff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-7.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-7.c
index 1b1be9dbb666..0f8dd25a18dc 100644
--- a/gcc/testsuite/c-c++-common/gomp/declare-mapper-7.c
+++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-7.c
@@ -1,4 +1,4 @@
-/* { dg-do compile { target c++ } } */
+/* { dg-do compile } */
 
 struct Q {
   int *arr1;
diff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-8.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-8.c
index 86ddb942072c..dadca282711c 100644
--- a/gcc/testsuite/c-c++-common/gomp/declare-mapper-8.c
+++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-8.c
@@ -1,4 +1,4 @@
-/* { dg-do compile { target c++ } } */
+/* { dg-do compile } */
 
 struct Q {
   int *arr1;
diff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-9.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-9.c
index 54e58426910e..b568c5a477f0 100644
--- a/gcc/testsuite/c-c++-common/gomp/declare-mapper-9.c
+++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-9.c
@@ -1,4 +1,4 @@
-/* { dg-do compile { target c++ } } */
+/* { dg-do compile } */
 
 int x = 5;
 
diff --git a/gcc/testsuite/gcc.dg/gomp/declare-mapper-10.c b/gcc/testsuite/gcc.dg/gomp/declare-mapper-10.c
new file mode 100644
index 000000000000..efc9c1369158
--- /dev/null
+++ b/gcc/testsuite/gcc.dg/gomp/declare-mapper-10.c
@@ -0,0 +1,61 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+// "omp declare mapper" support -- check expansion in gimple.
+
+#include <stdlib.h>
+
+struct S {
+  int *ptr;
+  int size;
+};
+
+#define N 64
+
+#pragma omp declare mapper (struct S w) map(w.size, w.ptr, w.ptr[:w.size])
+#pragma omp declare mapper (foo:struct S w) map(to:w.size, w.ptr) \
+					    map(w.ptr[:w.size])
+
+int main (int argc, char *argv[])
+{
+  struct S s;
+  s.ptr = (int *) malloc (sizeof (int) * N);
+  s.size = N;
+
+#pragma omp declare mapper (bar:struct S w) map(w.size, w.ptr, w.ptr[:w.size])
+
+#pragma omp target
+  {
+    for (int i = 0; i < N; i++)
+      s.ptr[i]++;
+  }
+
+#pragma omp target map(tofrom: s)
+  {
+    for (int i = 0; i < N; i++)
+      s.ptr[i]++;
+  }
+
+#pragma omp target map(mapper(default), tofrom: s)
+  {
+    for (int i = 0; i < N; i++)
+      s.ptr[i]++;
+  }
+
+#pragma omp target map(mapper(foo), alloc: s)
+  {
+    for (int i = 0; i < N; i++)
+      s.ptr[i]++;
+  }
+
+#pragma omp target map(mapper(bar), tofrom: s)
+  {
+    for (int i = 0; i < N; i++)
+      s.ptr[i]++;
+  }
+
+  return 0;
+}
+
+/* { dg-final { scan-tree-dump-times {map\(struct:s \[len: 2\]\) map\(tofrom:s\.ptr \[len: [0-9]+\]\) map\(tofrom:s\.size \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: _[0-9]+\]\) map\(attach:s\.ptr \[bias: 0\]\)} 4 "gimple" { target c++ } } } */
+/* { dg-final { scan-tree-dump-times {map\(struct:s \[len: 2\]\) map\(to:s\.ptr \[len: [0-9]+\]\) map\(to:s\.size \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: _[0-9]+\]\) map\(attach:s\.ptr \[bias: 0\]\)} 1 "gimple" { target c++ } } } */
diff --git a/gcc/testsuite/gcc.dg/gomp/declare-mapper-11.c b/gcc/testsuite/gcc.dg/gomp/declare-mapper-11.c
new file mode 100644
index 000000000000..927065e5ea63
--- /dev/null
+++ b/gcc/testsuite/gcc.dg/gomp/declare-mapper-11.c
@@ -0,0 +1,33 @@
+// { dg-do compile }
+
+// Error-checking tests for "omp declare mapper".
+
+typedef struct {
+  int *ptr;
+  int size;
+} S;
+
+typedef struct {
+  int z;
+} Z;
+
+int main (int argc, char *argv[])
+{
+#pragma omp declare mapper (S v) map(v.size, v.ptr[:v.size])
+/* { dg-error "previous '#pragma omp declare mapper'" "" { target c } .-1 } */
+
+  /* This one's a duplicate.  */
+#pragma omp declare mapper (default: S v) map (to: v.size) map (v)
+/* { dg-error "redeclaration of '<default>' '#pragma omp declare mapper' for type 'S'" "" { target c } .-1 } */
+
+  /* ...and this one doesn't use a "base language identifier" for the mapper
+     name.  */
+#pragma omp declare mapper (case: S v) map (to: v.size)
+/* { dg-error "expected identifier or 'default'" "" { target c } .-1 } */
+
+  /* A non-struct/class/union type isn't supposed to work.  */
+#pragma omp declare mapper (name:Z [5]foo) map (foo[0].z)
+/* { dg-error "'Z\\\[5\\\]' is not a struct or union type in '#pragma omp declare mapper'" "" { target c } .-1 } */
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-10.c b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-10.c
index b0fa40929fbc..ca5aef4d9d38 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-10.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-10.c
@@ -1,4 +1,4 @@
-/* { dg-do run { target c++ } } */
+/* { dg-do run } */
 
 #include <string.h>
 #include <stdlib.h>
diff --git a/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-11.c b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-11.c
index b509ddc412c5..942d6a5a6ad2 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-11.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-11.c
@@ -1,4 +1,4 @@
-/* { dg-do run { target c++ } } */
+/* { dg-do run } */
 
 #include <string.h>
 #include <stdlib.h>
diff --git a/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-12.c b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-12.c
index cf8919c22edf..cbedee51683f 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-12.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-12.c
@@ -1,4 +1,4 @@
-/* { dg-do run { target c++ } } */
+/* { dg-do run } */
 
 #include <string.h>
 #include <stdlib.h>
diff --git a/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-13.c b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-13.c
index 99b7eedad90f..c4784ebafdd5 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-13.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-13.c
@@ -1,4 +1,4 @@
-/* { dg-do run { target c++ } } */
+/* { dg-do run } */
 
 #include <assert.h>
 
diff --git a/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-14.c b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-14.c
index e7108da25fef..3e6027e30508 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-14.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-14.c
@@ -1,4 +1,4 @@
-/* { dg-do run { target c++ } } */
+/* { dg-do run } */
 
 #include <stdlib.h>
 #include <assert.h>
diff --git a/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-9.c b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-9.c
index 9f85df53998a..324d53567787 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-9.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-9.c
@@ -1,4 +1,4 @@
-/* { dg-do run { target c++ } } */
+/* { dg-do run } */
 
 #include <string.h>
 #include <stdlib.h>
-- 
2.41.0



More information about the Gcc-patches mailing list