[3/4] OpenMP: lvalue parsing for map clauses (C++)

Message ID 6c491905cd64b5b6a53d762d27afd9b875866323.1636974589.git.julian@codesourcery.com
State New
Headers
Series OpenMP: Parsing of lvalues for "map" clauses for C andjC++ |

Commit Message

Julian Brown Nov. 15, 2021, 11:18 a.m. UTC
  This patch changes parsing for OpenMP map clauses in C++ to use the
generic expression parser, hence adds support for parsing general
lvalues (as required by OpenMP 5.0+).  So far only a few new types of
expression are actually supported throughout compilation (including
everything in the testsuite of course, and newly-added tests), and we
attempt to reject unsupported expressions in order to avoid surprises
for the user.

The intention is to incrementally add support for further kinds of
lvalues on top of this patch.

OK?

Thanks,

Julian

2021-11-15  Julian Brown  <julian@codesourcery.com>

gcc/c-family/
	* c-omp.c (c_omp_decompose_attachable_address): Handle
	more types of expressions.

gcc/cp/
	* error.c (dump_expr): Handle OMP_ARRAY_SECTION.
	* parser.c (cp_parser_new): Initialize parser->omp_array_section_p.
	(cp_parser_postfix_open_square_expression): Support OMP_ARRAY_SECTION
	parsing.
	(cp_parser_omp_var_list_no_open): Remove ALLOW_DEREF parameter, add
	MAP_LVALUE in its place.  Supported generalised lvalue parsing for map
	clauses.
	(cp_parser_omp_var_list): Remove ALLOW_DEREF parameter, add MAP_LVALUE.
	Pass to cp_parser_omp_var_list_no_open.
	(cp_parser_oacc_data_clause, cp_parser_omp_all_clauses): Update calls
	to cp_parser_omp_var_list.
	* parser.h (cp_parser): Add omp_array_section_p field.
	* semantics.c (handle_omp_array_sections_1): Handle more types of map
	expression.
	(handle_omp_array_section): Handle non-DECL_P attachment points.
	(finish_omp_clauses): Check for supported types of expression.

gcc/
	* gimplify.c (build_struct_group): Handle reference-typed component
	accesses.  Fix support for non-DECL_P struct bases.
	(omp_build_struct_sibling_lists): Support length-two group for
	synthesized inner struct mapping.
	* tree-pretty-print.c (dump_generic_node): Support OMP_ARRAY_SECTION.
	* tree.def (OMP_ARRAY_SECTION): New tree code.

gcc/testsuite/
	* c-c++-common/gomp/map-6.c: Update expected output.
	* g++.dg/gomp/pr67522.C: Likewise.

libgomp/
	* testsuite/libgomp.c++/ind-base-1.C: New test.
	* testsuite/libgomp.c++/ind-base-2.C: New test.
	* testsuite/libgomp.c++/struct-ref-1.C: New test.
---
 gcc/c-family/c-omp.c                         |  25 ++-
 gcc/cp/error.c                               |   9 ++
 gcc/cp/parser.c                              | 141 ++++++++++++++--
 gcc/cp/parser.h                              |   3 +
 gcc/cp/semantics.c                           |  35 +++-
 gcc/gimplify.c                               |  37 ++++-
 gcc/testsuite/c-c++-common/gomp/map-6.c      |   4 +-
 gcc/testsuite/g++.dg/gomp/pr67522.C          |   2 +-
 gcc/tree-pretty-print.c                      |  14 ++
 gcc/tree.def                                 |   3 +
 libgomp/testsuite/libgomp.c++/ind-base-1.C   | 162 +++++++++++++++++++
 libgomp/testsuite/libgomp.c++/ind-base-2.C   |  49 ++++++
 libgomp/testsuite/libgomp.c++/struct-ref-1.C |  97 +++++++++++
 13 files changed, 553 insertions(+), 28 deletions(-)
 create mode 100644 libgomp/testsuite/libgomp.c++/ind-base-1.C
 create mode 100644 libgomp/testsuite/libgomp.c++/ind-base-2.C
 create mode 100644 libgomp/testsuite/libgomp.c++/struct-ref-1.C
  

Patch

diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c
index 9b85d47f78a..3834bd50cc9 100644
--- a/gcc/c-family/c-omp.c
+++ b/gcc/c-family/c-omp.c
@@ -3323,12 +3323,18 @@  c_omp_decompose_attachable_address (tree t, tree *virtbase)
 {
   *virtbase = t;
 
-  /* It's already a pointer.  Just use that.  */
-  if (POINTER_TYPE_P (TREE_TYPE (t)))
+  /* It's already a non-offset pointer.  Just use that.  */
+  if (POINTER_TYPE_P (TREE_TYPE (t))
+      && (DECL_P (t)
+	  || TREE_CODE (t) == COMPONENT_REF
+	  || TREE_CODE (t) == ARRAY_REF))
     return NULL_TREE;
 
   /* Otherwise, look for a base pointer deeper within the expression.  */
 
+  while (TREE_CODE (t) == COMPOUND_EXPR)
+    t = TREE_OPERAND (t, 1);
+
   while (TREE_CODE (t) == COMPONENT_REF
 	 && (TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF
 	     || TREE_CODE (TREE_OPERAND (t, 0)) == ARRAY_REF))
@@ -3338,9 +3344,24 @@  c_omp_decompose_attachable_address (tree t, tree *virtbase)
 	t = TREE_OPERAND (t, 0);
     }
 
+  if (TREE_CODE (t) == POINTER_PLUS_EXPR)
+    {
+      t = TREE_OPERAND (t, 0);
+      if (TREE_CODE (t) == SAVE_EXPR)
+	t = TREE_OPERAND (t, 0);
+    }
+
+  if (TREE_CODE (t) == INDIRECT_REF
+      && TREE_CODE (TREE_TYPE (TREE_OPERAND (t, 0))) == REFERENCE_TYPE)
+    t = TREE_OPERAND (t, 0);
 
   *virtbase = t;
 
+  /* If we have a pointer now (e.g. after we've stripped POINTER_PLUS_EXPR),
+     we have an offset pointer.  That's the attachment point.  */
+  if (POINTER_TYPE_P (TREE_TYPE (t)))
+    return t;
+
   if (TREE_CODE (t) != COMPONENT_REF)
     return NULL_TREE;
 
diff --git a/gcc/cp/error.c b/gcc/cp/error.c
index 012a4ecddf4..141c959d1bb 100644
--- a/gcc/cp/error.c
+++ b/gcc/cp/error.c
@@ -2415,6 +2415,15 @@  dump_expr (cxx_pretty_printer *pp, tree t, int flags)
       pp_cxx_right_bracket (pp);
       break;
 
+    case OMP_ARRAY_SECTION:
+      dump_expr (pp, TREE_OPERAND (t, 0), flags);
+      pp_cxx_left_bracket (pp);
+      dump_expr (pp, TREE_OPERAND (t, 1), flags);
+      pp_colon (pp);
+      dump_expr (pp, TREE_OPERAND (t, 2), flags);
+      pp_cxx_right_bracket (pp);
+      break;
+
     case UNARY_PLUS_EXPR:
       dump_unary_op (pp, "+", t, flags);
       break;
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 4f9e364aa95..2b2ad9b77d2 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -4241,6 +4241,9 @@  cp_parser_new (cp_lexer *lexer)
   parser->omp_declare_simd = NULL;
   parser->oacc_routine = NULL;
 
+  /* Allow array slice in expression.  */
+  parser->omp_array_section_p = false;
+
   /* Not declaring an implicit function template.  */
   parser->auto_is_implicit_function_template_parm_p = false;
   parser->fully_implicit_function_template_p = false;
@@ -7898,6 +7901,7 @@  cp_parser_postfix_open_square_expression (cp_parser *parser,
   tree index = NULL_TREE;
   location_t loc = cp_lexer_peek_token (parser->lexer)->location;
   bool saved_greater_than_is_operator_p;
+  bool saved_colon_corrects_to_scope_p;
 
   /* Consume the `[' token.  */
   cp_lexer_consume_token (parser->lexer);
@@ -7905,6 +7909,9 @@  cp_parser_postfix_open_square_expression (cp_parser *parser,
   saved_greater_than_is_operator_p = parser->greater_than_is_operator_p;
   parser->greater_than_is_operator_p = true;
 
+  saved_colon_corrects_to_scope_p = parser->colon_corrects_to_scope_p;
+  parser->colon_corrects_to_scope_p = false;
+
   /* Parse the index expression.  */
   /* ??? For offsetof, there is a question of what to allow here.  If
      offsetof is not being used in an integral constant expression context,
@@ -7915,7 +7922,8 @@  cp_parser_postfix_open_square_expression (cp_parser *parser,
      constant expressions here.  */
   if (for_offsetof)
     index = cp_parser_constant_expression (parser);
-  else
+  else if (!parser->omp_array_section_p
+	   || cp_lexer_next_token_is_not (parser->lexer, CPP_COLON))
     {
       if (cp_lexer_next_token_is (parser->lexer, CPP_OPEN_BRACE))
 	{
@@ -7932,6 +7940,32 @@  cp_parser_postfix_open_square_expression (cp_parser *parser,
 
   parser->greater_than_is_operator_p = saved_greater_than_is_operator_p;
 
+  if (parser->omp_array_section_p
+      && cp_lexer_next_token_is (parser->lexer, CPP_COLON))
+    {
+      cp_lexer_consume_token (parser->lexer);
+      tree length = NULL_TREE;
+      if (cp_lexer_next_token_is_not (parser->lexer, CPP_CLOSE_SQUARE))
+	length = cp_parser_expression (parser);
+
+      parser->colon_corrects_to_scope_p = saved_colon_corrects_to_scope_p;
+
+      if ((index && error_operand_p (index))
+	  || (length && error_operand_p (length)))
+	return error_mark_node;
+
+      cp_parser_require (parser, CPP_CLOSE_SQUARE, RT_CLOSE_SQUARE);
+
+      /* NOTE: We are reusing using the type of the whole array as the type of
+	 the array section here, which isn't necessarily entirely correct.
+	 Might need revisiting.  */
+      return build3_loc (input_location, OMP_ARRAY_SECTION,
+			 TREE_TYPE (postfix_expression),
+			 postfix_expression, index, length);
+    }
+
+  parser->colon_corrects_to_scope_p = saved_colon_corrects_to_scope_p;
+
   /* Look for the closing `]'.  */
   cp_parser_require (parser, CPP_CLOSE_SQUARE, RT_CLOSE_SQUARE);
 
@@ -36340,7 +36374,7 @@  struct omp_dim
 static tree
 cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
 				tree list, bool *colon,
-				bool allow_deref = false)
+				bool map_lvalue = false)
 {
   auto_vec<omp_dim> dims;
   bool array_section_p;
@@ -36351,12 +36385,95 @@  cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
       parser->colon_corrects_to_scope_p = false;
       *colon = false;
     }
+  begin_scope (sk_omp, NULL);
   while (1)
     {
       tree name, decl;
 
       if (kind == OMP_CLAUSE_DEPEND || kind == OMP_CLAUSE_AFFINITY)
 	cp_parser_parse_tentatively (parser);
+      else if (map_lvalue && kind == OMP_CLAUSE_MAP)
+	{
+	  auto s = make_temp_override (parser->omp_array_section_p, true);
+	  token = cp_lexer_peek_token (parser->lexer);
+	  location_t loc = token->location;
+	  decl = cp_parser_assignment_expression (parser);
+
+	  dims.truncate (0);
+	  if (TREE_CODE (decl) == OMP_ARRAY_SECTION)
+	    {
+	      while (TREE_CODE (decl) == OMP_ARRAY_SECTION)
+		{
+		  tree low_bound = TREE_OPERAND (decl, 1);
+		  tree length = TREE_OPERAND (decl, 2);
+		  dims.safe_push (omp_dim (low_bound, length, loc, false));
+		  decl = TREE_OPERAND (decl, 0);
+		}
+
+	      while (TREE_CODE (decl) == ARRAY_REF
+		     || TREE_CODE (decl) == INDIRECT_REF
+		     || TREE_CODE (decl) == COMPOUND_EXPR)
+		{
+		  if (REFERENCE_REF_P (decl))
+		    break;
+
+		  if (TREE_CODE (decl) == COMPOUND_EXPR)
+		    {
+		      decl = TREE_OPERAND (decl, 1);
+		      STRIP_NOPS (decl);
+		    }
+		  else if (TREE_CODE (decl) == INDIRECT_REF)
+		    {
+		      dims.safe_push (omp_dim (integer_zero_node,
+					       integer_one_node, loc, true));
+		      decl = TREE_OPERAND (decl, 0);
+		    }
+		  else  /* ARRAY_REF. */
+		    {
+		      tree index = TREE_OPERAND (decl, 1);
+		      dims.safe_push (omp_dim (index, integer_one_node, loc,
+					       true));
+		      decl = TREE_OPERAND (decl, 0);
+		    }
+		}
+
+	      /* Bare references have their own special handling, so remove
+		 the explicit dereference added by convert_from_reference.  */
+	      if (REFERENCE_REF_P (decl))
+		decl = TREE_OPERAND (decl, 0);
+
+	      for (int i = dims.length () - 1; i >= 0; i--)
+		decl = tree_cons (dims[i].low_bound, dims[i].length, decl);
+	    }
+	  else if (TREE_CODE (decl) == INDIRECT_REF)
+	    {
+	      bool ref_p = REFERENCE_REF_P (decl);
+
+	      /* Turn *foo into the representation previously used for
+		 foo[0].  */
+	      decl = TREE_OPERAND (decl, 0);
+	      STRIP_NOPS (decl);
+
+	      /* ...but don't add the [0:1] representation for references
+		 (because they have special handling elsewhere).  */
+	      if (!ref_p)
+		decl = tree_cons (integer_zero_node, integer_one_node, decl);
+	    }
+	  else if (TREE_CODE (decl) == ARRAY_REF)
+	    {
+	      tree idx = TREE_OPERAND (decl, 1);
+
+	      decl = TREE_OPERAND (decl, 0);
+	      STRIP_NOPS (decl);
+
+	      decl = tree_cons (idx, integer_one_node, decl);
+	    }
+	  else if (TREE_CODE (decl) == NON_LVALUE_EXPR
+		   || CONVERT_EXPR_P (decl))
+	    decl = TREE_OPERAND (decl, 0);
+
+	  goto build_clause;
+	}
       token = cp_lexer_peek_token (parser->lexer);
       if (kind != 0
 	  && cp_parser_is_keyword (token, RID_THIS))
@@ -36426,8 +36543,7 @@  cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
 	    case OMP_CLAUSE_TO:
 	    start_component_ref:
 	      while (cp_lexer_next_token_is (parser->lexer, CPP_DOT)
-		     || (allow_deref
-			 && cp_lexer_next_token_is (parser->lexer, CPP_DEREF)))
+		     || cp_lexer_next_token_is (parser->lexer, CPP_DEREF))
 		{
 		  cpp_ttype ttype
 		    = cp_lexer_next_token_is (parser->lexer, CPP_DOT)
@@ -36513,9 +36629,7 @@  cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
 		   || kind == OMP_CLAUSE_TO)
 		  && !array_section_p
 		  && (cp_lexer_next_token_is (parser->lexer, CPP_DOT)
-		      || (allow_deref
-			  && cp_lexer_next_token_is (parser->lexer,
-						     CPP_DEREF))))
+		      || cp_lexer_next_token_is (parser->lexer, CPP_DEREF)))
 		{
 		  for (unsigned i = 0; i < dims.length (); i++)
 		    {
@@ -36551,6 +36665,7 @@  cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
 		cp_parser_parse_definitely (parser);
 	    }
 
+	build_clause:
 	  tree u = build_omp_clause (token->location, kind);
 	  OMP_CLAUSE_DECL (u) = decl;
 	  OMP_CLAUSE_CHAIN (u) = list;
@@ -36572,6 +36687,7 @@  cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
     {
       *colon = true;
       cp_parser_require (parser, CPP_COLON, RT_COLON);
+      finish_scope ();
       return list;
     }
 
@@ -36592,6 +36708,7 @@  cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
 	goto get_comma;
     }
 
+  finish_scope ();
   return list;
 }
 
@@ -36600,11 +36717,11 @@  cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
 
 static tree
 cp_parser_omp_var_list (cp_parser *parser, enum omp_clause_code kind, tree list,
-			bool allow_deref = false)
+			bool map_lvalue = false)
 {
   if (cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
     return cp_parser_omp_var_list_no_open (parser, kind, list, NULL,
-					   allow_deref);
+					   map_lvalue);
   return list;
 }
 
@@ -36671,7 +36788,7 @@  cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
       gcc_unreachable ();
     }
   tree nl, c;
-  nl = cp_parser_omp_var_list (parser, OMP_CLAUSE_MAP, list, true);
+  nl = cp_parser_omp_var_list (parser, OMP_CLAUSE_MAP, list, false);
 
   for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
     OMP_CLAUSE_SET_MAP_KIND (c, kind);
@@ -40018,12 +40135,12 @@  cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask,
 					      clauses);
 	  else
 	    clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_TO, clauses,
-					      true);
+					      false);
 	  c_name = "to";
 	  break;
 	case PRAGMA_OMP_CLAUSE_FROM:
 	  clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_FROM, clauses,
-					    true);
+					    false);
 	  c_name = "from";
 	  break;
 	case PRAGMA_OMP_CLAUSE_UNIFORM:
diff --git a/gcc/cp/parser.h b/gcc/cp/parser.h
index 3669587cebd..bebf8e6b16a 100644
--- a/gcc/cp/parser.h
+++ b/gcc/cp/parser.h
@@ -404,6 +404,9 @@  struct GTY(()) cp_parser {
   /* TRUE if omp::directive or omp::sequence attributes may not appear.  */
   bool omp_attrs_forbidden_p;
 
+  /* TRUE if an OpenMP array section is allowed.  */
+  bool omp_array_section_p;
+
   /* Tracks the function's template parameter list when declaring a function
      using generic type parameters.  This is either a new chain in the case of a
      fully implicit function template or an extension of the function's existing
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 1499e565a61..b02787ccbe5 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -5046,7 +5046,9 @@  handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
       ret = t_insp.get_deref_toplevel ();
       if (TREE_CODE (t) == FIELD_DECL)
 	ret = finish_non_static_data_member (t, NULL_TREE, NULL_TREE);
-      else if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL)
+      else if (!VAR_P (t)
+	       && (ort == C_ORT_ACC || !EXPR_P (t))
+	       && TREE_CODE (t) != PARM_DECL)
 	{
 	  if (processing_template_decl && TREE_CODE (t) != OVERLOAD)
 	    return NULL_TREE;
@@ -5635,7 +5637,9 @@  handle_omp_array_sections (tree c, enum c_omp_region_type ort)
 	  bool reference_always_pointer = true;
 	  tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
 				      OMP_CLAUSE_MAP);
-	  if (TREE_CODE (t) == COMPONENT_REF)
+	  if (TREE_CODE (t) == COMPONENT_REF
+	      || (TREE_CODE (t) == POINTER_PLUS_EXPR
+		  && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF))
 	    {
 	      OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
 
@@ -5663,6 +5667,13 @@  handle_omp_array_sections (tree c, enum c_omp_region_type ort)
 		}
 	      OMP_CLAUSE_SET_MAP_KIND (c2, k);
 	    }
+	  else if (ort != C_ORT_ACC && attach_pt && !DECL_P (attach_pt))
+	    {
+	      if (TREE_CODE (TREE_TYPE (t)) != POINTER_TYPE)
+		return false;
+
+	      OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
+	    }
 	  else
 	    OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
 	  OMP_CLAUSE_MAP_IMPLICIT (c2) = OMP_CLAUSE_MAP_IMPLICIT (c);
@@ -7848,6 +7859,15 @@  finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		    {
 		      t = t_insp.analyze_components (false);
 
+		      if (!t_insp.map_supported_p ())
+			{
+			  sorry_at (OMP_CLAUSE_LOCATION (c),
+				    "unsupported map expression %qE",
+				    OMP_CLAUSE_DECL (c));
+			  remove = true;
+			  break;
+			}
+
 		      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 			  && OMP_CLAUSE_MAP_IMPLICIT (c)
 			  && (bitmap_bit_p (&map_head, DECL_UID (t))
@@ -7930,6 +7950,14 @@  finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	      OMP_CLAUSE_DECL (c) = t_insp.get_deref_toplevel ();
 	    if (type_dependent_expression_p (t_insp.get_deref_toplevel ()))
 	      break;
+	    if (!t_insp.map_supported_p ())
+	      {
+		sorry_at (OMP_CLAUSE_LOCATION (c),
+			  "unsupported map expression %qE",
+			  OMP_CLAUSE_DECL (c));
+		remove = true;
+		break;
+	      }
 	    if (t == error_mark_node)
 	      {
 		remove = true;
@@ -7959,7 +7987,8 @@  finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 		  && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
 		      || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER
-		      || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH))
+		      || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH
+		      || (ort != C_ORT_ACC && EXPR_P (t))))
 		break;
 	      if (DECL_P (t))
 		error_at (OMP_CLAUSE_LOCATION (c),
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 4563cce3fc6..8b0d0dbd084 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -10287,6 +10287,7 @@  build_struct_group (enum omp_region_type region_type, enum tree_code code,
     {
       tree l = build_omp_clause (OMP_CLAUSE_LOCATION (grp_end), OMP_CLAUSE_MAP);
       gomp_map_kind k = attach ? GOMP_MAP_FORCE_PRESENT : GOMP_MAP_STRUCT;
+      tree *tail_chain;
 
       OMP_CLAUSE_SET_MAP_KIND (l, k);
 
@@ -10315,9 +10316,13 @@  build_struct_group (enum omp_region_type region_type, enum tree_code code,
 	    {
 	      OMP_CLAUSE_CHAIN (extra_node) = *insert_node_pos;
 	      OMP_CLAUSE_CHAIN (alloc_node) = extra_node;
+	      tail_chain = &OMP_CLAUSE_CHAIN (extra_node);
 	    }
 	  else
-	    OMP_CLAUSE_CHAIN (alloc_node) = *insert_node_pos;
+	    {
+	      OMP_CLAUSE_CHAIN (alloc_node) = *insert_node_pos;
+	      tail_chain = &OMP_CLAUSE_CHAIN (alloc_node);
+	    }
 
 	  *insert_node_pos = l;
 	}
@@ -10325,6 +10330,7 @@  build_struct_group (enum omp_region_type region_type, enum tree_code code,
 	{
 	  gcc_assert (*grp_start_p == grp_end);
 	  grp_start_p = insert_node_after (l, grp_start_p);
+	  tail_chain = &OMP_CLAUSE_CHAIN (*grp_start_p);
 	}
 
       tree noind = strip_indirections (base);
@@ -10389,8 +10395,7 @@  build_struct_group (enum omp_region_type region_type, enum tree_code code,
 	     || TREE_CODE (sdecl) == POINTER_PLUS_EXPR)
 	sdecl = TREE_OPERAND (sdecl, 0);
 
-      if (DECL_P (sdecl)
-	  && POINTER_TYPE_P (TREE_TYPE (sdecl))
+      if (POINTER_TYPE_P (TREE_TYPE (sdecl))
 	  && (region_type & ORT_TARGET))
 	{
 	  tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (grp_end),
@@ -10404,8 +10409,12 @@  build_struct_group (enum omp_region_type region_type, enum tree_code code,
 		       && (TREE_CODE (TREE_TYPE (TREE_OPERAND
 						  (TREE_OPERAND (base, 0), 0)))
 			   == REFERENCE_TYPE))));
-	  enum gomp_map_kind mkind = base_ref ? GOMP_MAP_FIRSTPRIVATE_REFERENCE
-					      : GOMP_MAP_FIRSTPRIVATE_POINTER;
+	  enum gomp_map_kind mkind;
+	  if (DECL_P (sdecl))
+	    mkind = base_ref ? GOMP_MAP_FIRSTPRIVATE_REFERENCE
+			     : GOMP_MAP_FIRSTPRIVATE_POINTER;
+	  else
+	    mkind = GOMP_MAP_ATTACH;
 	  OMP_CLAUSE_SET_MAP_KIND (c2, mkind);
 	  OMP_CLAUSE_DECL (c2) = sdecl;
 	  tree baddr = build_fold_addr_expr (base);
@@ -10421,9 +10430,21 @@  build_struct_group (enum omp_region_type region_type, enum tree_code code,
 	  OMP_CLAUSE_SIZE (c2)
 	    = fold_build2_loc (OMP_CLAUSE_LOCATION (grp_end), MINUS_EXPR,
 			       ptrdiff_type_node, baddr, decladdr);
-	  /* Insert after struct node.  */
-	  OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (l);
-	  OMP_CLAUSE_CHAIN (l) = c2;
+	  if (mkind == GOMP_MAP_FIRSTPRIVATE_POINTER
+	      || mkind == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+	    {
+	      /* Insert after struct node.  */
+	      OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (l);
+	      OMP_CLAUSE_CHAIN (l) = c2;
+	    }
+	  else  /* GOMP_MAP_ATTACH.  */
+	    {
+	      /* Insert after struct group.  */
+	      OMP_CLAUSE_CHAIN (c2) = *tail_chain;
+	      *tail_chain = c2;
+	      if (*grp_start_p == grp_end)
+		return &OMP_CLAUSE_CHAIN (*tail_chain);
+	    }
 	}
 
       return NULL;
diff --git a/gcc/testsuite/c-c++-common/gomp/map-6.c b/gcc/testsuite/c-c++-common/gomp/map-6.c
index 6ee59714847..c749db845b0 100644
--- a/gcc/testsuite/c-c++-common/gomp/map-6.c
+++ b/gcc/testsuite/c-c++-common/gomp/map-6.c
@@ -20,12 +20,12 @@  foo (void)
   ;
 
   #pragma omp target map (close a) /* { dg-error "'close' undeclared" "" { target c } } */ 
-  /* { dg-error "'close' has not been declared" "" { target c++ } .-1 } */ 
+  /* { dg-error "'close' was not declared in this scope" "" { target c++ } .-1 } */ 
   /* { dg-error "expected '\\)' before 'a'" "" { target *-*-* } .-2 } */
   ;
 
   #pragma omp target map (always a) /* { dg-error "'always' undeclared" "" { target c } } */
-  /* { dg-error "'always' has not been declared" "" { target c++ } .-1 } */ 
+  /* { dg-error "'always' was not declared in this scope" "" { target c++ } .-1 } */ 
   /* { dg-error "expected '\\)' before 'a'" "" { target *-*-* } .-2 } */
   ;
 
diff --git a/gcc/testsuite/g++.dg/gomp/pr67522.C b/gcc/testsuite/g++.dg/gomp/pr67522.C
index da8cb74d1fa..4a901ba68c7 100644
--- a/gcc/testsuite/g++.dg/gomp/pr67522.C
+++ b/gcc/testsuite/g++.dg/gomp/pr67522.C
@@ -12,7 +12,7 @@  foo (void)
   for (int i = 0; i < 16; i++)
     ;
 
-  #pragma omp target map (S[0:10])		// { dg-error "is not a variable in" }
+  #pragma omp target map (S[0:10])		// { dg-error "expected primary-expression before '\\\[' token" }
   ;
 
   #pragma omp task depend (inout: S[0:10])	// { dg-error "is not a variable in" }
diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index b213bb9cec5..3ee27d5d25d 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -2510,6 +2510,20 @@  dump_generic_node (pretty_printer *pp, tree node, int spc, dump_flags_t flags,
 	}
       break;
 
+    case OMP_ARRAY_SECTION:
+      op0 = TREE_OPERAND (node, 0);
+      if (op_prio (op0) < op_prio (node))
+	pp_left_paren (pp);
+      dump_generic_node (pp, op0, spc, flags, false);
+      if (op_prio (op0) < op_prio (node))
+	pp_right_paren (pp);
+      pp_left_bracket (pp);
+      dump_generic_node (pp, TREE_OPERAND (node, 1), spc, flags, false);
+      pp_colon (pp);
+      dump_generic_node (pp, TREE_OPERAND (node, 2), spc, flags, false);
+      pp_right_bracket (pp);
+      break;
+
     case CONSTRUCTOR:
       {
 	unsigned HOST_WIDE_INT ix;
diff --git a/gcc/tree.def b/gcc/tree.def
index e27bc3e2b1f..9824840ec00 100644
--- a/gcc/tree.def
+++ b/gcc/tree.def
@@ -1304,6 +1304,9 @@  DEFTREECODE (OMP_ATOMIC_CAPTURE_NEW, "omp_atomic_capture_new", tcc_statement, 2)
 /* OpenMP clauses.  */
 DEFTREECODE (OMP_CLAUSE, "omp_clause", tcc_exceptional, 0)
 
+/* An OpenMP array section.  */
+DEFTREECODE (OMP_ARRAY_SECTION, "omp_array_section", tcc_expression, 3)
+
 /* TRANSACTION_EXPR tree code.
    Operand 0: BODY: contains body of the transaction.  */
 DEFTREECODE (TRANSACTION_EXPR, "transaction_expr", tcc_expression, 1)
diff --git a/libgomp/testsuite/libgomp.c++/ind-base-1.C b/libgomp/testsuite/libgomp.c++/ind-base-1.C
new file mode 100644
index 00000000000..4566854e60a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/ind-base-1.C
@@ -0,0 +1,162 @@ 
+// { dg-do run }
+// { dg-options "-fopenmp" }
+
+#include <cassert>
+
+struct S
+{
+  int x[10];
+};
+
+struct T
+{
+  struct S *s;
+};
+
+struct U
+{
+  struct T *t;
+};
+
+void
+foo_siblist (void)
+{
+  U *u = new U;
+  u->t = new T;
+  u->t->s = new S;
+  for (int i = 0; i < 10; i++)
+    u->t->s->x[i] = 0;
+#pragma omp target map(u->t, *(u->t), u->t->s, *u->t->s)
+  for (int i = 0; i < 10; i++)
+    u->t->s->x[i] = i * 3;
+  for (int i = 0; i < 10; i++)
+    assert (u->t->s->x[i] == i * 3);
+  delete u->t->s;
+  delete u->t;
+  delete u;
+}
+
+void
+foo (void)
+{
+  U *u = new U;
+  u->t = new T;
+  u->t->s = new S;
+  for (int i = 0; i < 10; i++)
+    u->t->s->x[i] = 0;
+#pragma omp target map(*u, u->t, *(u->t), u->t->s, *u->t->s)
+  for (int i = 0; i < 10; i++)
+    u->t->s->x[i] = i * 3;
+  for (int i = 0; i < 10; i++)
+    assert (u->t->s->x[i] == i * 3);
+  delete u->t->s;
+  delete u->t;
+  delete u;
+}
+
+void
+foo_tofrom (void)
+{
+  U *u = new U;
+  u->t = new T;
+  u->t->s = new S;
+  for (int i = 0; i < 10; i++)
+    u->t->s->x[i] = 0;
+#pragma omp target map(u, *u, u->t, *(u->t), u->t->s, *u->t->s)
+  for (int i = 0; i < 10; i++)
+    u->t->s->x[i] = i * 3;
+  for (int i = 0; i < 10; i++)
+    assert (u->t->s->x[i] == i * 3);
+  delete u->t->s;
+  delete u->t;
+  delete u;
+}
+
+void
+bar (void)
+{
+  U *u = new U;
+  U **up = &u;
+  u->t = new T;
+  u->t->s = new S;
+  for (int i = 0; i < 10; i++)
+    (*up)->t->s->x[i] = 0;
+#pragma omp target map(*up, (*up)->t, *(*up)->t, (*up)->t->s, *(*up)->t->s)
+  for (int i = 0; i < 10; i++)
+    (*up)->t->s->x[i] = i * 3;
+  for (int i = 0; i < 10; i++)
+    assert ((*up)->t->s->x[i] == i * 3);
+  delete u->t->s;
+  delete u->t;
+  delete u;
+}
+
+void
+bar_pp (void)
+{
+  U *u = new U;
+  U **up = &u;
+  u->t = new T;
+  u->t->s = new S;
+  for (int i = 0; i < 10; i++)
+    (*up)->t->s->x[i] = 0;
+#pragma omp target map(*up, **up, (*up)->t, *(*up)->t, (*up)->t->s, *(*up)->t->s)
+  for (int i = 0; i < 10; i++)
+    (*up)->t->s->x[i] = i * 3;
+  for (int i = 0; i < 10; i++)
+    assert ((*up)->t->s->x[i] == i * 3);
+  delete u->t->s;
+  delete u->t;
+  delete u;
+}
+
+void
+bar_tofrom (void)
+{
+  U *u = new U;
+  U **up = &u;
+  u->t = new T;
+  u->t->s = new S;
+  for (int i = 0; i < 10; i++)
+    (*up)->t->s->x[i] = 0;
+#pragma omp target map(*up, up, (*up)->t, *(*up)->t, (*up)->t->s, *(*up)->t->s)
+  for (int i = 0; i < 10; i++)
+    (*up)->t->s->x[i] = i * 3;
+  for (int i = 0; i < 10; i++)
+    assert ((*up)->t->s->x[i] == i * 3);
+  delete u->t->s;
+  delete u->t;
+  delete u;
+}
+
+void
+bar_tofrom_pp (void)
+{
+  U *u = new U;
+  U **up = &u;
+  u->t = new T;
+  u->t->s = new S;
+  for (int i = 0; i < 10; i++)
+    (*up)->t->s->x[i] = 0;
+#pragma omp target map(**up, *up, up, (*up)->t, *(*up)->t, (*up)->t->s, \
+		       *(*up)->t->s)
+  for (int i = 0; i < 10; i++)
+    (*up)->t->s->x[i] = i * 3;
+  for (int i = 0; i < 10; i++)
+    assert ((*up)->t->s->x[i] == i * 3);
+  delete u->t->s;
+  delete u->t;
+  delete u;
+}
+
+int main (int argc, char *argv[])
+{
+  foo_siblist ();
+  foo ();
+  foo_tofrom ();
+  bar ();
+  bar_pp ();
+  bar_tofrom ();
+  bar_tofrom_pp ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/ind-base-2.C b/libgomp/testsuite/libgomp.c++/ind-base-2.C
new file mode 100644
index 00000000000..706a1205c00
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/ind-base-2.C
@@ -0,0 +1,49 @@ 
+// { dg-do run }
+// { dg-options "-fopenmp" }
+
+#include <cassert>
+
+struct S
+{
+  int x[10];
+};
+
+struct T
+{
+  struct S ***s;
+};
+
+struct U
+{
+  struct T **t;
+};
+
+void
+foo (void)
+{
+  U *u = new U;
+  T *real_t = new T;
+  S *real_s = new S;
+  T **t_pp = &real_t;
+  S **s_pp = &real_s;
+  S ***s_ppp = &s_pp;
+  u->t = t_pp;
+  (*u->t)->s = s_ppp;
+  for (int i = 0; i < 10; i++)
+    (**((*u->t)->s))->x[i] = 0;
+#pragma omp target map(u->t, *u->t, (*u->t)->s, *(*u->t)->s, **(*u->t)->s, \
+		       (**(*u->t)->s)->x[0:10])
+  for (int i = 0; i < 10; i++)
+    (**((*u->t)->s))->x[i] = i * 3;
+  for (int i = 0; i < 10; i++)
+    assert ((**((*u->t)->s))->x[i] == i * 3);
+  delete real_s;
+  delete real_t;
+  delete u;
+}
+
+int main (int argc, char *argv[])
+{
+  foo ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/struct-ref-1.C b/libgomp/testsuite/libgomp.c++/struct-ref-1.C
new file mode 100644
index 00000000000..d3874650017
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/struct-ref-1.C
@@ -0,0 +1,97 @@ 
+// { dg-do run }
+// { dg-options "-fopenmp" }
+
+#include <cassert>
+
+struct S
+{
+  int x[10];
+};
+
+void
+foo (S *s, int x)
+{
+  S *&r = s;
+  for (int i = 0; i < x; i++)
+    s[0].x[i] = s[1].x[i] = 0;
+  #pragma omp target map (s, x)
+    ;
+  #pragma omp target map (s[0], x)
+  for (int i = 0; i < x; i++)
+    s[0].x[i] = i;
+  #pragma omp target map (s[1], x)
+  for (int i = 0; i < x; i++)
+    s[1].x[i] = i * 2;
+  for (int i = 0; i < x; i++)
+    {
+      assert (s[0].x[i] == i);
+      assert (s[1].x[i] == i * 2);
+      s[0].x[i] = 0;
+      s[1].x[i] = 0;
+    }
+  #pragma omp target map (r, x)
+    ;
+  #pragma omp target map (r[0], x)
+  for (int i = 0; i < x; i++)
+    r[0].x[i] = i;
+  #pragma omp target map (r[1], x)
+  for (int i = 0; i < x; i++)
+    r[1].x[i] = i * 2;
+  for (int i = 0; i < x; i++)
+    {
+      assert (r[0].x[i] == i);
+      assert (r[1].x[i] == i * 2);
+    }
+}
+
+template <int N>
+struct T
+{
+  int x[N];
+};
+
+template <int N>
+void
+bar (T<N> *t, int x)
+{
+  T<N> *&r = t;
+  for (int i = 0; i < x; i++)
+    t[0].x[i] = t[1].x[i] = 0;
+  #pragma omp target map (t, x)
+    ;
+  #pragma omp target map (t[0], x)
+  for (int i = 0; i < x; i++)
+    t[0].x[i] = i;
+  #pragma omp target map (t[1], x)
+  for (int i = 0; i < x; i++)
+    t[1].x[i] = i * 2;
+  for (int i = 0; i < x; i++)
+    {
+      assert (t[0].x[i] == i);
+      assert (t[1].x[i] == i * 2);
+      t[0].x[i] = 0;
+      t[1].x[i] = 0;
+    }
+  #pragma omp target map (r, x)
+    ;
+  #pragma omp target map (r[0], x)
+  for (int i = 0; i < x; i++)
+    r[0].x[i] = i;
+  #pragma omp target map (r[1], x)
+  for (int i = 0; i < x; i++)
+    r[1].x[i] = i * 2;
+  for (int i = 0; i < x; i++)
+    {
+      assert (r[0].x[i] == i);
+      assert (r[1].x[i] == i * 2);
+    }
+}
+
+int main (int argc, char *argv[])
+{
+  S s[2];
+  foo (s, 10);
+  T<10> t[2];
+  bar (t, 10);
+  return 0;
+}