diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index acbf20dcb58..1f27f35fdb9 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -12970,6 +12970,15 @@ c_parser_oacc_wait_list (c_parser *parser, location_t clause_loc, tree list)
    The optional ALLOW_DEREF argument is true if list items can use the deref
    (->) operator.  */
 
+struct omp_dim
+{
+  tree low_bound, length;
+  location_t loc;
+  bool no_colon;
+  omp_dim (tree lb, tree len, location_t lo, bool nc)
+  : low_bound (lb), length (len), loc (lo), no_colon (nc) {}
+};
+
 static tree
 c_parser_omp_variable_list (c_parser *parser,
 			    location_t clause_loc,
@@ -12982,6 +12991,7 @@ c_parser_omp_variable_list (c_parser *parser,
 
   while (1)
     {
+      auto_vec<omp_dim> dims;
       bool array_section_p = false;
       if (kind == OMP_CLAUSE_DEPEND || kind == OMP_CLAUSE_AFFINITY)
 	{
@@ -13101,6 +13111,7 @@ c_parser_omp_variable_list (c_parser *parser,
 	    case OMP_CLAUSE_MAP:
 	    case OMP_CLAUSE_FROM:
 	    case OMP_CLAUSE_TO:
+	    start_component_ref:
 	      while (c_parser_next_token_is (parser, CPP_DOT)
 		     || (allow_deref
 			 && c_parser_next_token_is (parser, CPP_DEREF)))
@@ -13128,9 +13139,13 @@ c_parser_omp_variable_list (c_parser *parser,
 	    case OMP_CLAUSE_REDUCTION:
 	    case OMP_CLAUSE_IN_REDUCTION:
 	    case OMP_CLAUSE_TASK_REDUCTION:
+	      array_section_p = false;
+	      dims.truncate (0);
 	      while (c_parser_next_token_is (parser, CPP_OPEN_SQUARE))
 		{
+		  location_t loc = UNKNOWN_LOCATION;
 		  tree low_bound = NULL_TREE, length = NULL_TREE;
+		  bool no_colon = false;
 
 		  c_parser_consume_token (parser);
 		  if (!c_parser_next_token_is (parser, CPP_COLON))
@@ -13141,9 +13156,13 @@ c_parser_omp_variable_list (c_parser *parser,
 		      expr = convert_lvalue_to_rvalue (expr_loc, expr,
 						       false, true);
 		      low_bound = expr.value;
+		      loc = expr_loc;
 		    }
 		  if (c_parser_next_token_is (parser, CPP_CLOSE_SQUARE))
-		    length = integer_one_node;
+		    {
+		      length = integer_one_node;
+		      no_colon = true;
+		    }
 		  else
 		    {
 		      /* Look for `:'.  */
@@ -13172,8 +13191,35 @@ c_parser_omp_variable_list (c_parser *parser,
 		      break;
 		    }
 
-		  t = tree_cons (low_bound, length, t);
+		  dims.safe_push (omp_dim (low_bound, length, loc, no_colon));
+		}
+
+	      if (t != error_mark_node)
+		{
+		  if ((kind == OMP_CLAUSE_MAP
+		       || kind == OMP_CLAUSE_FROM
+		       || kind == OMP_CLAUSE_TO)
+		      && !array_section_p
+		      && (c_parser_next_token_is (parser, CPP_DOT)
+			  || (allow_deref
+			      && c_parser_next_token_is (parser,
+							 CPP_DEREF))))
+		    {
+		      for (unsigned i = 0; i < dims.length (); i++)
+			{
+			  gcc_assert (dims[i].length == integer_one_node);
+			  t = build_array_ref (dims[i].loc,
+					       t, dims[i].low_bound);
+			}
+		      goto start_component_ref;
+		    }
+		  else
+		    {
+		      for (unsigned i = 0; i < dims.length (); i++)
+			t = tree_cons (dims[i].low_bound, dims[i].length, t);
+		    }
 		}
+
 	      if ((kind == OMP_CLAUSE_DEPEND || kind == OMP_CLAUSE_AFFINITY)
 		  && t != error_mark_node
 		  && parser->tokens_avail != 2)
@@ -16420,7 +16466,7 @@ c_parser_omp_clause_device_type (c_parser *parser, tree list)
 static tree
 c_parser_omp_clause_to (c_parser *parser, tree list)
 {
-  return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_TO, list);
+  return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_TO, list, true);
 }
 
 /* OpenMP 4.0:
@@ -16429,7 +16475,7 @@ c_parser_omp_clause_to (c_parser *parser, tree list)
 static tree
 c_parser_omp_clause_from (c_parser *parser, tree list)
 {
-  return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_FROM, list);
+  return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_FROM, list, true);
 }
 
 /* OpenMP 4.0:
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index c0ebb319aff..ee6362d4274 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -13217,6 +13217,18 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 		    t, omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
 	  return error_mark_node;
 	}
+      while (TREE_CODE (t) == INDIRECT_REF)
+	{
+	  t = TREE_OPERAND (t, 0);
+	  STRIP_NOPS (t);
+	  if (TREE_CODE (t) == POINTER_PLUS_EXPR)
+	    t = TREE_OPERAND (t, 0);
+	}
+      while (TREE_CODE (t) == COMPOUND_EXPR)
+	{
+	  t = TREE_OPERAND (t, 1);
+	  STRIP_NOPS (t);
+	}
       if (TREE_CODE (t) == COMPONENT_REF
 	  && (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 	      || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO
@@ -13238,10 +13250,14 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 		  return error_mark_node;
 		}
 	      t = TREE_OPERAND (t, 0);
-	      if (TREE_CODE (t) == MEM_REF)
+	      while (TREE_CODE (t) == MEM_REF
+		     || TREE_CODE (t) == INDIRECT_REF
+		     || TREE_CODE (t) == ARRAY_REF)
 		{
 		  t = TREE_OPERAND (t, 0);
 		  STRIP_NOPS (t);
+		  if (TREE_CODE (t) == POINTER_PLUS_EXPR)
+		    t = TREE_OPERAND (t, 0);
 		}
 	      if (ort == C_ORT_ACC && TREE_CODE (t) == MEM_REF)
 		{
@@ -13530,15 +13546,25 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 	  return error_mark_node;
 	}
       /* If there is a pointer type anywhere but in the very first
-	 array-section-subscript, the array section can't be contiguous.  */
+	 array-section-subscript, the array section could be non-contiguous.  */
       if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND
 	  && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_AFFINITY
 	  && TREE_CODE (TREE_CHAIN (t)) == TREE_LIST)
 	{
-	  error_at (OMP_CLAUSE_LOCATION (c),
-		    "array section is not contiguous in %qs clause",
-		    omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
-	  return error_mark_node;
+	  /* If any prior dimension has a non-one length, then deem this
+	     array section as non-contiguous.  */
+	  for (tree d = TREE_CHAIN (t); TREE_CODE (d) == TREE_LIST;
+	       d = TREE_CHAIN (d))
+	    {
+	      tree d_length = TREE_VALUE (d);
+	      if (d_length == NULL_TREE || !integer_onep (d_length))
+		{
+		  error_at (OMP_CLAUSE_LOCATION (c),
+			    "array section is not contiguous in %qs clause",
+			    omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+		  return error_mark_node;
+		}
+	    }
 	}
     }
   else
@@ -14888,13 +14914,20 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		  if (TREE_CODE (t) == COMPONENT_REF
 		      && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE)
 		    {
-		      while (TREE_CODE (t) == COMPONENT_REF)
-			t = TREE_OPERAND (t, 0);
-		      if (TREE_CODE (t) == MEM_REF)
+		      do
 			{
 			  t = TREE_OPERAND (t, 0);
-			  STRIP_NOPS (t);
+			  if (TREE_CODE (t) == MEM_REF
+			      || TREE_CODE (t) == INDIRECT_REF)
+			    {
+			      t = TREE_OPERAND (t, 0);
+			      STRIP_NOPS (t);
+			      if (TREE_CODE (t) == POINTER_PLUS_EXPR)
+				t = TREE_OPERAND (t, 0);
+			    }
 			}
+		      while (TREE_CODE (t) == COMPONENT_REF);
+
 		      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 			  && OMP_CLAUSE_MAP_IMPLICIT (c)
 			  && (bitmap_bit_p (&map_head, DECL_UID (t))
@@ -14961,14 +14994,32 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	       bias) to zero here, so it is not set erroneously to the pointer
 	       size later on in gimplify.c.  */
 	    OMP_CLAUSE_SIZE (c) = size_zero_node;
+	  while (TREE_CODE (t) == INDIRECT_REF
+		 || TREE_CODE (t) == ARRAY_REF)
+	    {
+	      t = TREE_OPERAND (t, 0);
+	      STRIP_NOPS (t);
+	      if (TREE_CODE (t) == POINTER_PLUS_EXPR)
+		t = TREE_OPERAND (t, 0);
+	    }
+	  while (TREE_CODE (t) == COMPOUND_EXPR)
+	    {
+	      t = TREE_OPERAND (t, 1);
+	      STRIP_NOPS (t);
+	    }
 	  indir_component_ref_p = false;
 	  if (TREE_CODE (t) == COMPONENT_REF
-	      && TREE_CODE (TREE_OPERAND (t, 0)) == MEM_REF)
+	      && (TREE_CODE (TREE_OPERAND (t, 0)) == MEM_REF
+		  || TREE_CODE (TREE_OPERAND (t, 0)) == INDIRECT_REF
+		  || TREE_CODE (TREE_OPERAND (t, 0)) == ARRAY_REF))
 	    {
 	      t = TREE_OPERAND (TREE_OPERAND (t, 0), 0);
 	      indir_component_ref_p = true;
 	      STRIP_NOPS (t);
+	      if (TREE_CODE (t) == POINTER_PLUS_EXPR)
+		t = TREE_OPERAND (t, 0);
 	    }
+
 	  if (TREE_CODE (t) == COMPONENT_REF
 	      && OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_)
 	    {
@@ -15004,7 +15055,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		      break;
 		    }
 		  t = TREE_OPERAND (t, 0);
-		  if (ort == C_ORT_ACC && TREE_CODE (t) == MEM_REF)
+		  if (TREE_CODE (t) == MEM_REF)
 		    {
 		      if (maybe_ne (mem_ref_offset (t), 0))
 			error_at (OMP_CLAUSE_LOCATION (c),
@@ -15013,6 +15064,15 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		      else
 			t = TREE_OPERAND (t, 0);
 		    }
+		  while (TREE_CODE (t) == MEM_REF
+			 || TREE_CODE (t) == INDIRECT_REF
+			 || TREE_CODE (t) == ARRAY_REF)
+		    {
+		      t = TREE_OPERAND (t, 0);
+		      STRIP_NOPS (t);
+		      if (TREE_CODE (t) == POINTER_PLUS_EXPR)
+			t = TREE_OPERAND (t, 0);
+		    }
 		}
 	      if (remove)
 		break;
@@ -15084,7 +15144,8 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 			    "%qD appears more than once in data clauses", t);
 		  remove = true;
 		}
-	      else if (bitmap_bit_p (&map_head, DECL_UID (t)))
+	      else if (bitmap_bit_p (&map_head, DECL_UID (t))
+		       && !bitmap_bit_p (&map_field_head, DECL_UID (t)))
 		{
 		  if (ort == C_ORT_ACC)
 		    error_at (OMP_CLAUSE_LOCATION (c),
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 8c27ea12013..fceea16a976 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -36323,11 +36323,22 @@ check_no_duplicate_clause (tree clauses, enum omp_clause_code code,
    The optional ALLOW_DEREF argument is true if list items can use the deref
    (->) operator.  */
 
+struct omp_dim
+{
+  tree low_bound, length;
+  location_t loc;
+  bool no_colon;
+  omp_dim (tree lb, tree len, location_t lo, bool nc)
+    : low_bound (lb), length (len), loc (lo), no_colon (nc) {}
+};
+
 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)
 {
+  auto_vec<omp_dim> dims;
+  bool array_section_p;
   cp_token *token;
   bool saved_colon_corrects_to_scope_p = parser->colon_corrects_to_scope_p;
   if (colon)
@@ -36408,6 +36419,7 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
 	    case OMP_CLAUSE_MAP:
 	    case OMP_CLAUSE_FROM:
 	    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)))
@@ -36431,14 +36443,19 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
 	    case OMP_CLAUSE_REDUCTION:
 	    case OMP_CLAUSE_IN_REDUCTION:
 	    case OMP_CLAUSE_TASK_REDUCTION:
+	      array_section_p = false;
+	      dims.truncate (0);
 	      while (cp_lexer_next_token_is (parser->lexer, CPP_OPEN_SQUARE))
 		{
+		  location_t loc = UNKNOWN_LOCATION;
 		  tree low_bound = NULL_TREE, length = NULL_TREE;
+		  bool no_colon = false;
 
 		  parser->colon_corrects_to_scope_p = false;
 		  cp_lexer_consume_token (parser->lexer);
 		  if (!cp_lexer_next_token_is (parser->lexer, CPP_COLON))
 		    {
+		      loc = cp_lexer_peek_token (parser->lexer)->location;
 		      low_bound = cp_parser_expression (parser);
 		      /* Later handling is not prepared to see through these.  */
 		      gcc_checking_assert (!location_wrapper_p (low_bound));
@@ -36447,7 +36464,10 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
 		    parser->colon_corrects_to_scope_p
 		      = saved_colon_corrects_to_scope_p;
 		  if (cp_lexer_next_token_is (parser->lexer, CPP_CLOSE_SQUARE))
-		    length = integer_one_node;
+		    {
+		      length = integer_one_node;
+		      no_colon = true;
+		    }
 		  else
 		    {
 		      /* Look for `:'.  */
@@ -36460,6 +36480,8 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
 			}
 		      if (kind == OMP_CLAUSE_DEPEND || kind == OMP_CLAUSE_AFFINITY)
 			cp_parser_commit_to_tentative_parse (parser);
+		      else
+			array_section_p = true;
 		      if (!cp_lexer_next_token_is (parser->lexer,
 						   CPP_CLOSE_SQUARE))
 			{
@@ -36478,8 +36500,32 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
 		      goto skip_comma;
 		    }
 
-		  decl = tree_cons (low_bound, length, decl);
+		  dims.safe_push (omp_dim (low_bound, length, loc, no_colon));
+		}
+
+	      if ((kind == OMP_CLAUSE_MAP
+		   || kind == OMP_CLAUSE_FROM
+		   || 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))))
+		{
+		  for (unsigned i = 0; i < dims.length (); i++)
+		    {
+		      gcc_assert (dims[i].length == integer_one_node);
+		      decl = build_array_ref (dims[i].loc,
+					      decl, dims[i].low_bound);
+		    }
+		  goto start_component_ref;
 		}
+	      else
+		{
+		  for (unsigned i = 0; i < dims.length (); i++)
+		    decl = tree_cons (dims[i].low_bound, dims[i].length, decl);
+		}
+
 	      break;
 	    default:
 	      break;
@@ -39981,11 +40027,13 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask,
 	    clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_TO_DECLARE,
 					      clauses);
 	  else
-	    clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_TO, clauses);
+	    clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_TO, clauses,
+					      true);
 	  c_name = "to";
 	  break;
 	case PRAGMA_OMP_CLAUSE_FROM:
-	  clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_FROM, clauses);
+	  clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_FROM, clauses,
+					    true);
 	  c_name = "from";
 	  break;
 	case PRAGMA_OMP_CLAUSE_UNIFORM:
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index c64b45c0cee..898fcea6bb6 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -5030,6 +5030,18 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 	  && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
 	t = TREE_OPERAND (t, 0);
       ret = t;
+      while (TREE_CODE (t) == INDIRECT_REF)
+	{
+	  t = TREE_OPERAND (t, 0);
+	  STRIP_NOPS (t);
+	  if (TREE_CODE (t) == POINTER_PLUS_EXPR)
+	    t = TREE_OPERAND (t, 0);
+	}
+      while (TREE_CODE (t) == COMPOUND_EXPR)
+	{
+	  t = TREE_OPERAND (t, 1);
+	  STRIP_NOPS (t);
+	}
       if (TREE_CODE (t) == COMPONENT_REF
 	  && (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 	      || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO
@@ -5054,10 +5066,14 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 		  return error_mark_node;
 		}
 	      t = TREE_OPERAND (t, 0);
-	      if (TREE_CODE (t) == INDIRECT_REF)
+	      while (TREE_CODE (t) == MEM_REF
+		     || TREE_CODE (t) == INDIRECT_REF
+		     || TREE_CODE (t) == ARRAY_REF)
 		{
 		  t = TREE_OPERAND (t, 0);
 		  STRIP_NOPS (t);
+		  if (TREE_CODE (t) == POINTER_PLUS_EXPR)
+		    t = TREE_OPERAND (t, 0);
 		}
 	    }
 	  if (REFERENCE_REF_P (t))
@@ -5341,15 +5357,25 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 	  return error_mark_node;
 	}
       /* If there is a pointer type anywhere but in the very first
-	 array-section-subscript, the array section can't be contiguous.  */
+	 array-section-subscript, the array section could be non-contiguous.  */
       if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_AFFINITY
 	  && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND
 	  && TREE_CODE (TREE_CHAIN (t)) == TREE_LIST)
 	{
-	  error_at (OMP_CLAUSE_LOCATION (c),
-		    "array section is not contiguous in %qs clause",
-		    omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
-	  return error_mark_node;
+	  /* If any prior dimension has a non-one length, then deem this
+	     array section as non-contiguous.  */
+	  for (tree d = TREE_CHAIN (t); TREE_CODE (d) == TREE_LIST;
+	       d = TREE_CHAIN (d))
+	    {
+	      tree d_length = TREE_VALUE (d);
+	      if (d_length == NULL_TREE || !integer_onep (d_length))
+		{
+		  error_at (OMP_CLAUSE_LOCATION (c),
+			    "array section is not contiguous in %qs clause",
+			    omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+		  return error_mark_node;
+		}
+	    }
 	}
     }
   else
@@ -5619,16 +5645,37 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
 	    default:
 	      break;
 	    }
+	  bool reference_always_pointer = true;
 	  tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
 				      OMP_CLAUSE_MAP);
 	  if (TREE_CODE (t) == COMPONENT_REF)
-	    OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
+	    {
+	      OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
+
+	      if ((ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP
+		  && TYPE_REF_P (TREE_TYPE (t)))
+		{
+		  if (TREE_CODE (TREE_TYPE (TREE_TYPE (t))) == ARRAY_TYPE)
+		    OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER);
+		  else
+		    t = convert_from_reference (t);
+
+		  reference_always_pointer = false;
+		}
+	    }
 	  else if (REFERENCE_REF_P (t)
 		   && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
 	    {
-	      t = TREE_OPERAND (t, 0);
-	      gomp_map_kind k = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH
-						   : GOMP_MAP_ALWAYS_POINTER;
+	      gomp_map_kind k;
+	      if ((ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP
+		  && TREE_CODE (TREE_TYPE (t)) == POINTER_TYPE)
+		k = GOMP_MAP_ATTACH_DETACH;
+	      else
+		{
+		  t = TREE_OPERAND (t, 0);
+		  k = (ort == C_ORT_ACC
+		       ? GOMP_MAP_ATTACH_DETACH : GOMP_MAP_ALWAYS_POINTER);
+		}
 	      OMP_CLAUSE_SET_MAP_KIND (c2, k);
 	    }
 	  else
@@ -5652,8 +5699,10 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
 	  OMP_CLAUSE_SIZE (c2) = t;
 	  OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
 	  OMP_CLAUSE_CHAIN (c) = c2;
+
 	  ptr = OMP_CLAUSE_DECL (c2);
-	  if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
+	  if (reference_always_pointer
+	      && OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
 	      && TYPE_REF_P (TREE_TYPE (ptr))
 	      && INDIRECT_TYPE_P (TREE_TYPE (TREE_TYPE (ptr))))
 	    {
@@ -7855,15 +7904,22 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		  if (TREE_CODE (t) == COMPONENT_REF
 		      && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE)
 		    {
-		      while (TREE_CODE (t) == COMPONENT_REF)
-			t = TREE_OPERAND (t, 0);
-		      if (REFERENCE_REF_P (t))
-			t = TREE_OPERAND (t, 0);
-		      if (TREE_CODE (t) == INDIRECT_REF)
+		      do
 			{
 			  t = TREE_OPERAND (t, 0);
-			  STRIP_NOPS (t);
+			  if (REFERENCE_REF_P (t))
+			    t = TREE_OPERAND (t, 0);
+			  if (TREE_CODE (t) == MEM_REF
+			      || TREE_CODE (t) == INDIRECT_REF)
+			    {
+			      t = TREE_OPERAND (t, 0);
+			      STRIP_NOPS (t);
+			      if (TREE_CODE (t) == POINTER_PLUS_EXPR)
+				t = TREE_OPERAND (t, 0);
+			    }
 			}
+		      while (TREE_CODE (t) == COMPONENT_REF);
+
 		      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 			  && OMP_CLAUSE_MAP_IMPLICIT (c)
 			  && (bitmap_bit_p (&map_head, DECL_UID (t))
@@ -7934,15 +7990,33 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	      && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
 	    {
 	      t = TREE_OPERAND (t, 0);
-	      OMP_CLAUSE_DECL (c) = t;
+	      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH_DETACH)
+		OMP_CLAUSE_DECL (c) = t;
+	    }
+	  while (TREE_CODE (t) == INDIRECT_REF
+		 || TREE_CODE (t) == ARRAY_REF)
+	    {
+	      t = TREE_OPERAND (t, 0);
+	      STRIP_NOPS (t);
+	      if (TREE_CODE (t) == POINTER_PLUS_EXPR)
+		t = TREE_OPERAND (t, 0);
+	    }
+	  while (TREE_CODE (t) == COMPOUND_EXPR)
+	    {
+	      t = TREE_OPERAND (t, 1);
+	      STRIP_NOPS (t);
 	    }
 	  indir_component_ref_p = false;
 	  if (TREE_CODE (t) == COMPONENT_REF
-	      && TREE_CODE (TREE_OPERAND (t, 0)) == INDIRECT_REF)
+	      && (TREE_CODE (TREE_OPERAND (t, 0)) == INDIRECT_REF
+		  || TREE_CODE (TREE_OPERAND (t, 0)) == ARRAY_REF))
 	    {
 	      t = TREE_OPERAND (TREE_OPERAND (t, 0), 0);
 	      indir_component_ref_p = true;
 	      STRIP_NOPS (t);
+	      if (TREE_CODE (t) == POINTER_PLUS_EXPR)
+		t = TREE_OPERAND (t, 0);
 	    }
 	  if (TREE_CODE (t) == COMPONENT_REF
 	      && OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_)
@@ -7977,6 +8051,24 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		      break;
 		    }
 		  t = TREE_OPERAND (t, 0);
+		  if (TREE_CODE (t) == MEM_REF)
+		    {
+		      if (maybe_ne (mem_ref_offset (t), 0))
+			error_at (OMP_CLAUSE_LOCATION (c),
+				  "cannot dereference %qE in %qs clause", t,
+				  omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+		      else
+			t = TREE_OPERAND (t, 0);
+		    }
+		  while (TREE_CODE (t) == MEM_REF
+			 || TREE_CODE (t) == INDIRECT_REF
+			 || TREE_CODE (t) == ARRAY_REF)
+		    {
+		      t = TREE_OPERAND (t, 0);
+		      STRIP_NOPS (t);
+		      if (TREE_CODE (t) == POINTER_PLUS_EXPR)
+			t = TREE_OPERAND (t, 0);
+		    }
 		}
 	      if (remove)
 		break;
@@ -8074,7 +8166,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 			    "%qD appears more than once in data clauses", t);
 		  remove = true;
 		}
-	      else if (bitmap_bit_p (&map_head, DECL_UID (t)))
+	      else if (bitmap_bit_p (&map_head, DECL_UID (t))
+		       && !bitmap_bit_p (&map_field_head, DECL_UID (t)))
 		{
 		  if (ort == C_ORT_ACC)
 		    error_at (OMP_CLAUSE_LOCATION (c),
@@ -8121,8 +8214,13 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	  else
 	    {
 	      bitmap_set_bit (&map_head, DECL_UID (t));
-	      if (t != OMP_CLAUSE_DECL (c)
-		  && TREE_CODE (OMP_CLAUSE_DECL (c)) == COMPONENT_REF)
+
+	      tree decl = OMP_CLAUSE_DECL (c);
+	      if (t != decl
+		  && (TREE_CODE (decl) == COMPONENT_REF
+		      || (INDIRECT_REF_P (decl)
+			  && TREE_CODE (TREE_OPERAND (decl, 0)) == COMPONENT_REF
+			  && TYPE_REF_P (TREE_TYPE (TREE_OPERAND (decl, 0))))))
 		bitmap_set_bit (&map_field_head, DECL_UID (t));
 	    }
 	handle_map_references:
@@ -8151,7 +8249,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		  tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
 					      OMP_CLAUSE_MAP);
 		  if (TREE_CODE (t) == COMPONENT_REF)
-		    OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
+		    OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER);
 		  else
 		    OMP_CLAUSE_SET_MAP_KIND (c2,
 					     GOMP_MAP_FIRSTPRIVATE_REFERENCE);
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index b86c7cf9833..a762d715f4c 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -2460,6 +2460,9 @@ gfc_trans_omp_array_section (stmtblock_t *block, gfc_omp_namelist *n,
 			       TREE_TYPE (TREE_TYPE (decl)),
 			       decl, offset, NULL_TREE, NULL_TREE);
 	  OMP_CLAUSE_DECL (node) = offset;
+
+	  if (ptr_kind == GOMP_MAP_ALWAYS_POINTER)
+	    return;
 	}
       else
 	{
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index ed46fe3c461..9d37766f853 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -8666,7 +8666,7 @@ insert_struct_comp_map (enum tree_code code, tree c, tree struct_node,
 
 static tree
 extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp,
-			 poly_offset_int *poffsetp)
+			 poly_offset_int *poffsetp, tree *offsetp)
 {
   tree offset;
   poly_int64 bitsize, bitpos;
@@ -8713,10 +8713,11 @@ extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp,
       && TREE_CODE (TREE_TYPE (TREE_OPERAND (base, 0))) == REFERENCE_TYPE)
     base = TREE_OPERAND (base, 0);
 
-  gcc_assert (offset == NULL_TREE || poly_int_tree_p (offset));
-
-  if (offset)
-    poffset = wi::to_poly_offset (offset);
+  if (offset && poly_int_tree_p (offset))
+    {
+      poffset = wi::to_poly_offset (offset);
+      offset = NULL_TREE;
+    }
   else
     poffset = 0;
 
@@ -8725,6 +8726,7 @@ extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp,
 
   *bitposp = bitpos;
   *poffsetp = poffset;
+  *offsetp = offset;
 
   /* Set *BASE_REF if BASE was a dereferenced reference variable.  */
   if (base_ref && orig_base != base)
@@ -8738,12 +8740,22 @@ extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp,
 static bool
 is_or_contains_p (tree expr, tree base_ptr)
 {
-  while (expr != base_ptr)
-    if (TREE_CODE (base_ptr) == COMPONENT_REF)
-      base_ptr = TREE_OPERAND (base_ptr, 0);
-    else
-      break;
-  return expr == base_ptr;
+  if ((TREE_CODE (expr) == INDIRECT_REF && TREE_CODE (base_ptr) == MEM_REF)
+      || (TREE_CODE (expr) == MEM_REF && TREE_CODE (base_ptr) == INDIRECT_REF))
+    return operand_equal_p (TREE_OPERAND (expr, 0),
+			    TREE_OPERAND (base_ptr, 0));
+  while (!operand_equal_p (expr, base_ptr))
+    {
+      if (TREE_CODE (base_ptr) == COMPOUND_EXPR)
+	base_ptr = TREE_OPERAND (base_ptr, 1);
+      if (TREE_CODE (base_ptr) == COMPONENT_REF
+	  || TREE_CODE (base_ptr) == POINTER_PLUS_EXPR
+	  || TREE_CODE (base_ptr) == SAVE_EXPR)
+	base_ptr = TREE_OPERAND (base_ptr, 0);
+      else
+	break;
+    }
+  return operand_equal_p (expr, base_ptr);
 }
 
 /* Implement OpenMP 5.x map ordering rules for target directives. There are
@@ -8823,21 +8835,107 @@ omp_target_reorder_clauses (tree *list_p)
 	    tree base_ptr = TREE_OPERAND (decl, 0);
 	    STRIP_TYPE_NOPS (base_ptr);
 	    for (unsigned int j = i + 1; j < atf.length (); j++)
-	      {
-		tree *cp2 = atf[j];
-		tree decl2 = OMP_CLAUSE_DECL (*cp2);
-		if (is_or_contains_p (decl2, base_ptr))
-		  {
-		    /* Move *cp2 to before *cp.  */
-		    tree c = *cp2;
-		    *cp2 = OMP_CLAUSE_CHAIN (c);
-		    OMP_CLAUSE_CHAIN (c) = *cp;
-		    *cp = c;
-		    atf[j] = NULL;
+	      if (atf[j])
+		{
+		  tree *cp2 = atf[j];
+		  tree decl2 = OMP_CLAUSE_DECL (*cp2);
+
+		  decl2 = OMP_CLAUSE_DECL (*cp2);
+		  if (is_or_contains_p (decl2, base_ptr))
+		    {
+		      /* Move *cp2 to before *cp.  */
+		      tree c = *cp2;
+		      *cp2 = OMP_CLAUSE_CHAIN (c);
+		      OMP_CLAUSE_CHAIN (c) = *cp;
+		      *cp = c;
+
+		      if (*cp2 != NULL_TREE
+			  && OMP_CLAUSE_CODE (*cp2) == OMP_CLAUSE_MAP
+			  && OMP_CLAUSE_MAP_KIND (*cp2) == GOMP_MAP_ALWAYS_POINTER)
+			{
+			  tree c2 = *cp2;
+			  *cp2 = OMP_CLAUSE_CHAIN (c2);
+			  OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
+			  OMP_CLAUSE_CHAIN (c) = c2;
+			}
+
+		      atf[j] = NULL;
 		  }
-	      }
+		}
 	  }
       }
+
+  /* For attach_detach map clauses, if there is another map that maps the
+     attached/detached pointer, make sure that map is ordered before the
+     attach_detach.  */
+  atf.truncate (0);
+  for (tree *cp = list_p; *cp; cp = &OMP_CLAUSE_CHAIN (*cp))
+    if (OMP_CLAUSE_CODE (*cp) == OMP_CLAUSE_MAP)
+      {
+	/* Collect alloc, to, from, to/from clauses, and
+	   always_pointer/attach_detach clauses.  */
+	gomp_map_kind k = OMP_CLAUSE_MAP_KIND (*cp);
+	if (k == GOMP_MAP_ALLOC
+	    || k == GOMP_MAP_TO
+	    || k == GOMP_MAP_FROM
+	    || k == GOMP_MAP_TOFROM
+	    || k == GOMP_MAP_ALWAYS_TO
+	    || k == GOMP_MAP_ALWAYS_FROM
+	    || k == GOMP_MAP_ALWAYS_TOFROM
+	    || k == GOMP_MAP_ATTACH_DETACH
+	    || k == GOMP_MAP_ALWAYS_POINTER)
+	  atf.safe_push (cp);
+      }
+
+  for (unsigned int i = 0; i < atf.length (); i++)
+    if (atf[i])
+      {
+	tree *cp = atf[i];
+	tree ptr = OMP_CLAUSE_DECL (*cp);
+	STRIP_TYPE_NOPS (ptr);
+	if (OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_ATTACH_DETACH)
+	  for (unsigned int j = i + 1; j < atf.length (); j++)
+	    {
+	      tree *cp2 = atf[j];
+	      tree decl2 = OMP_CLAUSE_DECL (*cp2);
+	      if (OMP_CLAUSE_MAP_KIND (*cp2) != GOMP_MAP_ATTACH_DETACH
+		  && OMP_CLAUSE_MAP_KIND (*cp2) != GOMP_MAP_ALWAYS_POINTER
+		  && is_or_contains_p (decl2, ptr))
+		{
+		  /* Move *cp2 to before *cp.  */
+		  tree c = *cp2;
+		  *cp2 = OMP_CLAUSE_CHAIN (c);
+		  OMP_CLAUSE_CHAIN (c) = *cp;
+		  *cp = c;
+		  atf[j] = NULL;
+
+		  /* If decl2 is of the form '*decl2_opnd0', and followed by an
+		     ALWAYS_POINTER or ATTACH_DETACH of 'decl2_opnd0', move the
+		     pointer operation along with *cp2. This can happen for C++
+		     reference sequences.  */
+		  if (j + 1 < atf.length ()
+		      && (TREE_CODE (decl2) == INDIRECT_REF
+			  || TREE_CODE (decl2) == MEM_REF))
+		    {
+		      tree *cp3 = atf[j + 1];
+		      tree decl3 = OMP_CLAUSE_DECL (*cp3);
+		      tree decl2_opnd0 = TREE_OPERAND (decl2, 0);
+		      if ((OMP_CLAUSE_MAP_KIND (*cp3) == GOMP_MAP_ALWAYS_POINTER
+			   || OMP_CLAUSE_MAP_KIND (*cp3) == GOMP_MAP_ATTACH_DETACH)
+			  && operand_equal_p (decl3, decl2_opnd0))
+			{
+			  /* Also move *cp3 to before *cp.  */
+			  c = *cp3;
+			  *cp2 = OMP_CLAUSE_CHAIN (c);
+			  OMP_CLAUSE_CHAIN (c) = *cp;
+			  *cp = c;
+			  atf[j + 1] = NULL;
+			  j += 1;
+			}
+		    }
+		}
+	    }
+      }
 }
 
 /* DECL is supposed to have lastprivate semantics in the outer contexts
@@ -8929,6 +9027,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
   struct gimplify_omp_ctx *ctx, *outer_ctx;
   tree c;
   hash_map<tree_operand_hash, tree> *struct_map_to_clause = NULL;
+  hash_map<tree_operand_hash, tree *> *struct_seen_clause = NULL;
   hash_set<tree> *struct_deref_set = NULL;
   tree *prev_list_p = NULL, *orig_list_p = list_p;
   int handled_depend_iterators = -1;
@@ -9404,6 +9503,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		}
 	      bool indir_p = false;
 	      bool component_ref_p = false;
+	      tree indir_base = NULL_TREE;
 	      tree orig_decl = decl;
 	      tree decl_ref = NULL_TREE;
 	      if ((region_type & (ORT_ACC | ORT_TARGET | ORT_TARGET_DATA)) != 0
@@ -9422,6 +9522,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 			      == POINTER_TYPE))
 			{
 			  indir_p = true;
+			  indir_base = decl;
 			  decl = TREE_OPERAND (decl, 0);
 			  STRIP_NOPS (decl);
 			}
@@ -9468,7 +9569,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 			       != GOMP_MAP_POINTER)
 			   || OMP_CLAUSE_DECL (next_clause) != decl)
 		      && (!struct_deref_set
-			  || !struct_deref_set->contains (decl)))
+			  || !struct_deref_set->contains (decl))
+		      && (!struct_map_to_clause
+			  || !struct_map_to_clause->get (indir_base)))
 		    {
 		      if (!struct_deref_set)
 			struct_deref_set = new hash_set<tree> ();
@@ -9512,7 +9615,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	      if ((DECL_P (decl)
 		   || (component_ref_p
 		       && (INDIRECT_REF_P (decl)
-			   || TREE_CODE (decl) == MEM_REF)))
+			   || TREE_CODE (decl) == MEM_REF
+			   || TREE_CODE (decl) == ARRAY_REF)))
 		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET
 		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH
 		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_DETACH
@@ -9547,7 +9651,15 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 			  remove = true;
 			  break;
 			}
-		      if (OMP_CLAUSE_CHAIN (*prev_list_p) != c)
+
+		      /* The below prev_list_p based error recovery code is
+			 currently no longer valid for OpenMP.  */
+		      if (code != OMP_TARGET
+			  && code != OMP_TARGET_DATA
+			  && code != OMP_TARGET_UPDATE
+			  && code != OMP_TARGET_ENTER_DATA
+			  && code != OMP_TARGET_EXIT_DATA
+			  && OMP_CLAUSE_CHAIN (*prev_list_p) != c)
 			{
 			  tree ch = OMP_CLAUSE_CHAIN (*prev_list_p);
 			  if (ch == NULL_TREE || OMP_CLAUSE_CHAIN (ch) != c)
@@ -9560,13 +9672,15 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 
 		  poly_offset_int offset1;
 		  poly_int64 bitpos1;
+		  tree tree_offset1;
 		  tree base_ref;
 
 		  tree base
 		    = extract_base_bit_offset (OMP_CLAUSE_DECL (c), &base_ref,
-					       &bitpos1, &offset1);
+					       &bitpos1, &offset1,
+					       &tree_offset1);
 
-		  gcc_assert (base == decl);
+		  bool do_map_struct = (base == decl && !tree_offset1);
 
 		  splay_tree_node n
 		    = (DECL_P (decl)
@@ -9598,6 +9712,32 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		      OMP_CLAUSE_SET_MAP_KIND (c, k);
 		      has_attachments = true;
 		    }
+
+		  /* We currently don't handle non-constant offset accesses wrt to
+		     GOMP_MAP_STRUCT elements.  */
+		  if (!do_map_struct)
+		    goto skip_map_struct;
+
+		  /* Nor for attach_detach for OpenMP.  */
+		  if ((code == OMP_TARGET
+		       || code == OMP_TARGET_DATA
+		       || code == OMP_TARGET_UPDATE
+		       || code == OMP_TARGET_ENTER_DATA
+		       || code == OMP_TARGET_EXIT_DATA)
+		      && attach_detach)
+		    {
+		      if (DECL_P (decl))
+			{
+			  if (struct_seen_clause == NULL)
+			    struct_seen_clause
+			      = new hash_map<tree_operand_hash, tree *>;
+			  if (!struct_seen_clause->get (decl))
+			    struct_seen_clause->put (decl, list_p);
+			}
+
+		      goto skip_map_struct;
+		    }
+
 		  if ((DECL_P (decl)
 		       && (n == NULL || (n->value & GOVD_MAP) == 0))
 		      || (!DECL_P (decl)
@@ -9637,9 +9777,14 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		      struct_map_to_clause->put (decl, l);
 		      if (ptr || attach_detach)
 			{
-			  insert_struct_comp_map (code, c, l, *prev_list_p,
+			  tree **sc = (struct_seen_clause
+				       ? struct_seen_clause->get (decl)
+				       : NULL);
+			  tree *insert_node_pos = sc ? *sc : prev_list_p;
+
+			  insert_struct_comp_map (code, c, l, *insert_node_pos,
 						  NULL);
-			  *prev_list_p = l;
+			  *insert_node_pos = l;
 			  prev_list_p = NULL;
 			}
 		      else
@@ -9725,9 +9870,11 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 			    tree sc_decl = OMP_CLAUSE_DECL (*sc);
 			    poly_offset_int offsetn;
 			    poly_int64 bitposn;
+			    tree tree_offsetn;
 			    tree base
 			      = extract_base_bit_offset (sc_decl, NULL,
-							 &bitposn, &offsetn);
+							 &bitposn, &offsetn,
+							 &tree_offsetn);
 			    if (base != decl)
 			      break;
 			    if (scp)
@@ -9815,16 +9962,21 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 			  continue;
 			}
 		    }
+		skip_map_struct:
+		  ;
 		}
 	      else if ((code == OACC_ENTER_DATA
 			|| code == OACC_EXIT_DATA
 			|| code == OACC_DATA
 			|| code == OACC_PARALLEL
 			|| code == OACC_KERNELS
-			|| code == OACC_SERIAL)
+			|| code == OACC_SERIAL
+			|| code == OMP_TARGET_ENTER_DATA
+			|| code == OMP_TARGET_EXIT_DATA)
 		       && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
 		{
-		  gomp_map_kind k = (code == OACC_EXIT_DATA
+		  gomp_map_kind k = ((code == OACC_EXIT_DATA
+				      || code == OMP_TARGET_EXIT_DATA)
 				     ? GOMP_MAP_DETACH : GOMP_MAP_ATTACH);
 		  OMP_CLAUSE_SET_MAP_KIND (c, k);
 		}
@@ -10656,6 +10808,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 
   ctx->clauses = *orig_list_p;
   gimplify_omp_ctxp = ctx;
+  if (struct_seen_clause)
+    delete struct_seen_clause;
   if (struct_map_to_clause)
     delete struct_map_to_clause;
   if (struct_deref_set)
diff --git a/gcc/testsuite/c-c++-common/goacc/deep-copy-arrayofstruct.c b/gcc/testsuite/c-c++-common/goacc/deep-copy-arrayofstruct.c
index d411bcfa8e7..4247607b61c 100644
--- a/gcc/testsuite/c-c++-common/goacc/deep-copy-arrayofstruct.c
+++ b/gcc/testsuite/c-c++-common/goacc/deep-copy-arrayofstruct.c
@@ -37,13 +37,12 @@ int main(int argc, char* argv[])
     {
       int j, k;
       for (k = 0; k < S; k++)
-#pragma acc parallel loop copy(m[k].a[0:N]) /* { dg-error "expected .\\\). before .\\\.. token" } */
+#pragma acc parallel loop copy(m[k].a[0:N])
         for (j = 0; j < N; j++)
           m[k].a[j]++;
 
       for (k = 0; k < S; k++)
-#pragma acc parallel loop copy(m[k].b[0:N], m[k].c[5:N-10]) /* { dg-error "expected .\\\). before .\\\.. token" } */
-	/* { dg-error ".m. appears more than once in data clauses" "" { target c++ } .-1 } */
+#pragma acc parallel loop copy(m[k].b[0:N], m[k].c[5:N-10])
 	for (j = 0; j < N; j++)
 	  {
 	    m[k].b[j]++;
diff --git a/gcc/testsuite/c-c++-common/gomp/target-enter-data-1.c b/gcc/testsuite/c-c++-common/gomp/target-enter-data-1.c
new file mode 100644
index 00000000000..ce766d29e2d
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-enter-data-1.c
@@ -0,0 +1,24 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fopenmp -fdump-tree-gimple" } */
+
+struct bar
+{
+  int num_vectors;
+  double *vectors;
+};
+
+struct foo
+{
+  int num_vectors;
+  struct bar *bars;
+  double **vectors;
+};
+
+void func (struct foo *f, int n, int m)
+{
+  #pragma omp target enter data map (to: f->vectors[m][:n])
+  #pragma omp target enter data map (to: f->bars[n].vectors[:m])
+  #pragma omp target enter data map (to: f->bars[n].vectors[:f->bars[n].num_vectors])
+}
+
+/* { dg-final { scan-tree-dump-times "map\\(to:\\*_\[0-9\]+ \\\[len: _\[0-9\]+\\\]\\) map\\(attach:\[^-\]+->vectors \\\[bias: \[^\]\]+\\\]\\)" 3 "gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/target-implicit-map-2.c b/gcc/testsuite/c-c++-common/gomp/target-implicit-map-2.c
new file mode 100644
index 00000000000..3aa1a8fc55e
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-implicit-map-2.c
@@ -0,0 +1,52 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+#include <stdlib.h>
+
+#define N 10
+
+struct S
+{
+  int a, b;
+  int *ptr;
+  int c, d;
+};
+
+int
+main (void)
+{
+  struct S a;
+  a.ptr = (int *) malloc (sizeof (int) * N);
+
+  for (int i = 0; i < N; i++)
+    a.ptr[i] = 0;
+
+  #pragma omp target enter data map(to: a.ptr, a.ptr[:N])
+
+  #pragma omp target
+  for (int i = 0; i < N; i++)
+    a.ptr[i] += 1;
+
+  #pragma omp target update from(a.ptr[:N])
+
+  for (int i = 0; i < N; i++)
+    if (a.ptr[i] != 1)
+      abort ();
+
+  #pragma omp target map(a.ptr[:N])
+  for (int i = 0; i < N; i++)
+    a.ptr[i] += 1;
+
+  #pragma omp target update from(a.ptr[:N])
+
+  for (int i = 0; i < N; i++)
+    if (a.ptr[i] != 2)
+      abort ();
+
+  #pragma omp target exit data map(from:a.ptr, a.ptr[:N])
+
+  return 0;
+}
+
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(tofrom:a \[len: [0-9]+\]\[implicit\]\)} "gimple" } } */
+
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(tofrom:a \[len: [0-9]+\]\[implicit\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(attach:a\.ptr \[bias: 0\]\)} "gimple" } } */
diff --git a/libgomp/target.c b/libgomp/target.c
index bb31b1991d1..917da807d9c 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -581,11 +581,30 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep,
 	 address/length adjustment is a TODO.  */
       assert (!implicit_subset);
 
-      gomp_copy_host2dev (devicep, aq,
-			  (void *) (oldn->tgt->tgt_start + oldn->tgt_offset
-				    + newn->host_start - oldn->host_start),
-			  (void *) newn->host_start,
-			  newn->host_end - newn->host_start, false, cbuf);
+      if (oldn->aux && oldn->aux->attach_count)
+	{
+	  /* We have to be careful not to overwrite still attached pointers
+	     during the copyback to host.  */
+	  uintptr_t addr = newn->host_start;
+	  while (addr < newn->host_end)
+	    {
+	      size_t i = (addr - oldn->host_start) / sizeof (void *);
+	      if (oldn->aux->attach_count[i] == 0)
+		gomp_copy_host2dev (devicep, aq,
+				    (void *) (oldn->tgt->tgt_start
+					      + oldn->tgt_offset
+					      + addr - oldn->host_start),
+				    (void *) addr,
+				    sizeof (void *), false, cbuf);
+	      addr += sizeof (void *);
+	    }
+	}
+      else
+	gomp_copy_host2dev (devicep, aq,
+			    (void *) (oldn->tgt->tgt_start + oldn->tgt_offset
+				      + newn->host_start - oldn->host_start),
+			    (void *) newn->host_start,
+			    newn->host_end - newn->host_start, false, cbuf);
     }
 
   gomp_increment_refcount (oldn, refcount_set);
@@ -2011,17 +2030,45 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs,
 			    (void *) n->host_end);
 	      }
 
-
-	    void *hostaddr = (void *) cur_node.host_start;
-	    void *devaddr = (void *) (n->tgt->tgt_start + n->tgt_offset
-				      + cur_node.host_start - n->host_start);
-	    size_t size = cur_node.host_end - cur_node.host_start;
-
-	    if (GOMP_MAP_COPY_TO_P (kind & typemask))
-	      gomp_copy_host2dev (devicep, NULL, devaddr, hostaddr, size,
-				  false, NULL);
-	    if (GOMP_MAP_COPY_FROM_P (kind & typemask))
-	      gomp_copy_dev2host (devicep, NULL, hostaddr, devaddr, size);
+	    if (n->aux && n->aux->attach_count)
+	      {
+		uintptr_t addr = cur_node.host_start;
+		while (addr < cur_node.host_end)
+		  {
+		    /* We have to be careful not to overwrite still attached
+		       pointers during host<->device updates.  */
+		    size_t i = (addr - cur_node.host_start) / sizeof (void *);
+		    if (n->aux->attach_count[i] == 0)
+		      {
+			void *devaddr = (void *) (n->tgt->tgt_start
+						  + n->tgt_offset
+						  + addr - n->host_start);
+			if (GOMP_MAP_COPY_TO_P (kind & typemask))
+			  gomp_copy_host2dev (devicep, NULL,
+					      devaddr, (void *) addr,
+					      sizeof (void *), false, NULL);
+			if (GOMP_MAP_COPY_FROM_P (kind & typemask))
+			  gomp_copy_dev2host (devicep, NULL,
+					      (void *) addr, devaddr,
+					      sizeof (void *));
+		      }
+		    addr += sizeof (void *);
+		  }
+	      }
+	    else
+	      {
+		void *hostaddr = (void *) cur_node.host_start;
+		void *devaddr = (void *) (n->tgt->tgt_start + n->tgt_offset
+					  + cur_node.host_start
+					  - n->host_start);
+		size_t size = cur_node.host_end - cur_node.host_start;
+
+		if (GOMP_MAP_COPY_TO_P (kind & typemask))
+		  gomp_copy_host2dev (devicep, NULL, devaddr, hostaddr, size,
+				      false, NULL);
+		if (GOMP_MAP_COPY_FROM_P (kind & typemask))
+		  gomp_copy_dev2host (devicep, NULL, hostaddr, devaddr, size);
+	      }
 	  }
       }
   gomp_mutex_unlock (&devicep->lock);
@@ -2915,11 +2962,31 @@ gomp_exit_data (struct gomp_device_descr *devicep, size_t mapnum,
 
 	  if ((kind == GOMP_MAP_FROM && do_copy)
 	      || kind == GOMP_MAP_ALWAYS_FROM)
-	    gomp_copy_dev2host (devicep, NULL, (void *) cur_node.host_start,
-				(void *) (k->tgt->tgt_start + k->tgt_offset
-					  + cur_node.host_start
-					  - k->host_start),
-				cur_node.host_end - cur_node.host_start);
+	    {
+	      if (k->aux && k->aux->attach_count)
+		{
+		  /* We have to be careful not to overwrite still attached
+		     pointers during the copyback to host.  */
+		  uintptr_t addr = k->host_start;
+		  while (addr < k->host_end)
+		    {
+		      size_t i = (addr - k->host_start) / sizeof (void *);
+		      if (k->aux->attach_count[i] == 0)
+			gomp_copy_dev2host (devicep, NULL, (void *) addr,
+					    (void *) (k->tgt->tgt_start
+						      + k->tgt_offset
+						      + addr - k->host_start),
+					    sizeof (void *));
+		      addr += sizeof (void *);
+		    }
+		}
+	      else
+		gomp_copy_dev2host (devicep, NULL, (void *) cur_node.host_start,
+				    (void *) (k->tgt->tgt_start + k->tgt_offset
+					      + cur_node.host_start
+					      - k->host_start),
+				    cur_node.host_end - cur_node.host_start);
+	    }
 
 	  /* Structure elements lists are removed altogether at once, which
 	     may cause immediate deallocation of the target_mem_desc, causing
diff --git a/libgomp/testsuite/libgomp.c++/target-11.C b/libgomp/testsuite/libgomp.c++/target-11.C
index fe99603351d..87c2980b4b5 100644
--- a/libgomp/testsuite/libgomp.c++/target-11.C
+++ b/libgomp/testsuite/libgomp.c++/target-11.C
@@ -23,9 +23,11 @@ foo ()
   e = c + 18;
   D s = { a, b + 2, { 0, a + 16, 0, d }, c + 3, e };
   int err = 0;
-  #pragma omp target map (to:s.v.b[0:z + 7], s.template u[z + 1:z + 4]) \
-		     map (tofrom:s.s[3:3], s. template v. template d[z + 1:z + 3]) \
-		     map (from: s.w[z:4], s.x[1:3], err) private (i)
+  #pragma omp target map (to: s.v.b, s.v.b[0:z + 7])			\
+		     map (s.template u, s.template u[z + 1:z + 4])	\
+		     map (tofrom: s.s, s.s[3:3])			\
+		     map (tofrom: s. template v. template d[z + 1:z + 3])\
+		     map (from: s.w, s.w[z:4], s.x, s.x[1:3], err) private (i)
   {
     err = 0;
     for (i = 0; i < 7; i++)
@@ -80,9 +82,9 @@ main ()
   e = c + 18;
   S s = { a, b + 2, { 0, a + 16, 0, d }, c + 3, e };
   int err = 0;
-  #pragma omp target map (to:s.v.b[0:z + 7], s.u[z + 1:z + 4]) \
-		     map (tofrom:s.s[3:3], s.v.d[z + 1:z + 3]) \
-		     map (from: s.w[z:4], s.x[1:3], err) private (i)
+  #pragma omp target map (to: s.v.b, s.v.b[0:z + 7], s.u, s.u[z + 1:z + 4]) \
+		     map (tofrom: s.s, s.s[3:3], s.v.d[z + 1:z + 3])		\
+		     map (from: s.w, s.w[z:4], s.x, s.x[1:3], err) private (i)
   {
     err = 0;
     for (i = 0; i < 7; i++)
diff --git a/libgomp/testsuite/libgomp.c++/target-12.C b/libgomp/testsuite/libgomp.c++/target-12.C
index 3b4ed57df68..480e479c262 100644
--- a/libgomp/testsuite/libgomp.c++/target-12.C
+++ b/libgomp/testsuite/libgomp.c++/target-12.C
@@ -53,7 +53,7 @@ main ()
   int u[10] = { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9 }, err = 0;
   S s = { 9, u + 3, { 10, 11, 12, 13, 14 } };
   int *v = u + 4;
-  #pragma omp target enter data map (to: s.s, s.u[0:5]) map (alloc: s.v[1:3])
+  #pragma omp target enter data map (to: s.s, s.u, s.u[0:5]) map (alloc: s.v[1:3])
   s.s++;
   u[3]++;
   s.v[1]++;
diff --git a/libgomp/testsuite/libgomp.c++/target-15.C b/libgomp/testsuite/libgomp.c++/target-15.C
index 4b320c31229..53626b2547e 100644
--- a/libgomp/testsuite/libgomp.c++/target-15.C
+++ b/libgomp/testsuite/libgomp.c++/target-15.C
@@ -14,7 +14,7 @@ foo (S s)
     d = id;
 
   int err;
-  #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(to: sep) map(from: err)
+  #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(to: sep) map(from: err)
   {
     err = s.a != 11 || s.b[0] != 12 || s.b[1] != 13;
     err |= s.c[1] != 15 || s.c[2] != 16 || s.d[-2] != 18 || s.d[-1] != 19 || s.d[0] != 20;
@@ -48,7 +48,7 @@ foo (S s)
 	  || omp_target_is_present (&s.h, d)
 	  || omp_target_is_present (&s.h[2], d)))
     abort ();
-  #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
   {
     if (!omp_target_is_present (&s.a, d)
 	|| !omp_target_is_present (s.b, d)
@@ -61,8 +61,8 @@ foo (S s)
 	|| !omp_target_is_present (&s.h, d)
 	|| !omp_target_is_present (&s.h[2], d))
       abort ();
-    #pragma omp target update to(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
-    #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
+    #pragma omp target update to(s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
+    #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(from: err)
     {
       err = s.a != 50 || s.b[0] != 49 || s.b[1] != 48;
       err |= s.c[1] != 47 || s.c[2] != 46 || s.d[-2] != 45 || s.d[-1] != 44 || s.d[0] != 43;
@@ -73,7 +73,7 @@ foo (S s)
       s.e = 25; s.f[0] = 26; s.f[1] = 27; s.g[1] = 28; s.g[2] = 29;
       s.h[2] = 30; s.h[3] = 31; s.h[4] = 32;
     }
-    #pragma omp target update from(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+    #pragma omp target update from(s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
   }
   if (sep
       && (omp_target_is_present (&s.a, d)
@@ -97,7 +97,7 @@ foo (S s)
   s.c[1] = 36; s.c[2] = 37; s.d[-2] = 38; s.d[-1] = 39; s.d[0] = 40;
   s.e = 41; s.f[0] = 42; s.f[1] = 43; s.g[1] = 44; s.g[2] = 45;
   s.h[2] = 46; s.h[3] = 47; s.h[4] = 48;
-  #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
   if (!omp_target_is_present (&s.a, d)
       || !omp_target_is_present (s.b, d)
       || !omp_target_is_present (&s.c[1], d)
@@ -109,8 +109,8 @@ foo (S s)
       || !omp_target_is_present (&s.h, d)
       || !omp_target_is_present (&s.h[2], d))
     abort ();
-  #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
-  #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
+  #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
+  #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(from: err)
   {
     err = s.a != 33 || s.b[0] != 34 || s.b[1] != 35;
     err |= s.c[1] != 36 || s.c[2] != 37 || s.d[-2] != 38 || s.d[-1] != 39 || s.d[0] != 40;
@@ -121,7 +121,7 @@ foo (S s)
     s.e = 31; s.f[0] = 40; s.f[1] = 39; s.g[1] = 38; s.g[2] = 37;
     s.h[2] = 36; s.h[3] = 35; s.h[4] = 34;
   }
-  #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
   if (!omp_target_is_present (&s.a, d)
       || !omp_target_is_present (s.b, d)
       || !omp_target_is_present (&s.c[1], d)
@@ -133,7 +133,7 @@ foo (S s)
       || !omp_target_is_present (&s.h, d)
       || !omp_target_is_present (&s.h[2], d))
     abort ();
-  #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
   if (sep
       && (omp_target_is_present (&s.a, d)
 	  || omp_target_is_present (s.b, d)
diff --git a/libgomp/testsuite/libgomp.c++/target-16.C b/libgomp/testsuite/libgomp.c++/target-16.C
index cd102d90594..b8be7cc922f 100644
--- a/libgomp/testsuite/libgomp.c++/target-16.C
+++ b/libgomp/testsuite/libgomp.c++/target-16.C
@@ -16,7 +16,7 @@ foo (S<C, I, L, UC, SH> s)
     d = id;
 
   int err;
-  #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(to: sep) map(from: err)
+  #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(to: sep) map(from: err)
   {
     err = s.a != 11 || s.b[0] != 12 || s.b[1] != 13;
     err |= s.c[1] != 15 || s.c[2] != 16 || s.d[-2] != 18 || s.d[-1] != 19 || s.d[0] != 20;
@@ -50,7 +50,7 @@ foo (S<C, I, L, UC, SH> s)
 	  || omp_target_is_present (&s.h, d)
 	  || omp_target_is_present (&s.h[2], d)))
     abort ();
-  #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
   {
     if (!omp_target_is_present (&s.a, d)
 	|| !omp_target_is_present (s.b, d)
@@ -63,8 +63,8 @@ foo (S<C, I, L, UC, SH> s)
 	|| !omp_target_is_present (&s.h, d)
 	|| !omp_target_is_present (&s.h[2], d))
       abort ();
-    #pragma omp target update to(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
-    #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
+    #pragma omp target update to(s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
+    #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(from: err)
     {
       err = s.a != 50 || s.b[0] != 49 || s.b[1] != 48;
       err |= s.c[1] != 47 || s.c[2] != 46 || s.d[-2] != 45 || s.d[-1] != 44 || s.d[0] != 43;
@@ -75,7 +75,7 @@ foo (S<C, I, L, UC, SH> s)
       s.e = 25; s.f[0] = 26; s.f[1] = 27; s.g[1] = 28; s.g[2] = 29;
       s.h[2] = 30; s.h[3] = 31; s.h[4] = 32;
     }
-    #pragma omp target update from(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+    #pragma omp target update from(s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
   }
   if (sep
       && (omp_target_is_present (&s.a, d)
@@ -99,7 +99,7 @@ foo (S<C, I, L, UC, SH> s)
   s.c[1] = 36; s.c[2] = 37; s.d[-2] = 38; s.d[-1] = 39; s.d[0] = 40;
   s.e = 41; s.f[0] = 42; s.f[1] = 43; s.g[1] = 44; s.g[2] = 45;
   s.h[2] = 46; s.h[3] = 47; s.h[4] = 48;
-  #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
   if (!omp_target_is_present (&s.a, d)
       || !omp_target_is_present (s.b, d)
       || !omp_target_is_present (&s.c[1], d)
@@ -111,8 +111,8 @@ foo (S<C, I, L, UC, SH> s)
       || !omp_target_is_present (&s.h, d)
       || !omp_target_is_present (&s.h[2], d))
     abort ();
-  #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
-  #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
+  #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
+  #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(from: err)
   {
     err = s.a != 33 || s.b[0] != 34 || s.b[1] != 35;
     err |= s.c[1] != 36 || s.c[2] != 37 || s.d[-2] != 38 || s.d[-1] != 39 || s.d[0] != 40;
@@ -123,7 +123,7 @@ foo (S<C, I, L, UC, SH> s)
     s.e = 31; s.f[0] = 40; s.f[1] = 39; s.g[1] = 38; s.g[2] = 37;
     s.h[2] = 36; s.h[3] = 35; s.h[4] = 34;
   }
-  #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
   if (!omp_target_is_present (&s.a, d)
       || !omp_target_is_present (s.b, d)
       || !omp_target_is_present (&s.c[1], d)
@@ -135,7 +135,7 @@ foo (S<C, I, L, UC, SH> s)
       || !omp_target_is_present (&s.h, d)
       || !omp_target_is_present (&s.h[2], d))
     abort ();
-  #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
   if (sep
       && (omp_target_is_present (&s.a, d)
 	  || omp_target_is_present (s.b, d)
diff --git a/libgomp/testsuite/libgomp.c++/target-17.C b/libgomp/testsuite/libgomp.c++/target-17.C
index d81ff19a411..f97476aafc4 100644
--- a/libgomp/testsuite/libgomp.c++/target-17.C
+++ b/libgomp/testsuite/libgomp.c++/target-17.C
@@ -16,7 +16,7 @@ foo (S<C, I, L, UCR, CAR, SH, IPR> s)
     d = id;
 
   int err;
-  #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(to: sep) map(from: err)
+  #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(to: sep) map(from: err)
   {
     err = s.a != 11 || s.b[0] != 12 || s.b[1] != 13;
     err |= s.c[1] != 15 || s.c[2] != 16 || s.d[-2] != 18 || s.d[-1] != 19 || s.d[0] != 20;
@@ -50,7 +50,7 @@ foo (S<C, I, L, UCR, CAR, SH, IPR> s)
 	  || omp_target_is_present (&s.h, d)
 	  || omp_target_is_present (&s.h[2], d)))
     abort ();
-  #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
   {
     if (!omp_target_is_present (&s.a, d)
 	|| !omp_target_is_present (s.b, d)
@@ -63,8 +63,8 @@ foo (S<C, I, L, UCR, CAR, SH, IPR> s)
 	|| !omp_target_is_present (&s.h, d)
 	|| !omp_target_is_present (&s.h[2], d))
       abort ();
-    #pragma omp target update to(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
-    #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
+    #pragma omp target update to(s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
+    #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(from: err)
     {
       err = s.a != 50 || s.b[0] != 49 || s.b[1] != 48;
       err |= s.c[1] != 47 || s.c[2] != 46 || s.d[-2] != 45 || s.d[-1] != 44 || s.d[0] != 43;
@@ -75,7 +75,7 @@ foo (S<C, I, L, UCR, CAR, SH, IPR> s)
       s.e = 25; s.f[0] = 26; s.f[1] = 27; s.g[1] = 28; s.g[2] = 29;
       s.h[2] = 30; s.h[3] = 31; s.h[4] = 32;
     }
-    #pragma omp target update from(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+    #pragma omp target update from(s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
   }
   if (sep
       && (omp_target_is_present (&s.a, d)
@@ -99,7 +99,7 @@ foo (S<C, I, L, UCR, CAR, SH, IPR> s)
   s.c[1] = 36; s.c[2] = 37; s.d[-2] = 38; s.d[-1] = 39; s.d[0] = 40;
   s.e = 41; s.f[0] = 42; s.f[1] = 43; s.g[1] = 44; s.g[2] = 45;
   s.h[2] = 46; s.h[3] = 47; s.h[4] = 48;
-  #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
   if (!omp_target_is_present (&s.a, d)
       || !omp_target_is_present (s.b, d)
       || !omp_target_is_present (&s.c[1], d)
@@ -111,8 +111,8 @@ foo (S<C, I, L, UCR, CAR, SH, IPR> s)
       || !omp_target_is_present (&s.h, d)
       || !omp_target_is_present (&s.h[2], d))
     abort ();
-  #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
-  #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
+  #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
+  #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(from: err)
   {
     err = s.a != 33 || s.b[0] != 34 || s.b[1] != 35;
     err |= s.c[1] != 36 || s.c[2] != 37 || s.d[-2] != 38 || s.d[-1] != 39 || s.d[0] != 40;
@@ -123,7 +123,7 @@ foo (S<C, I, L, UCR, CAR, SH, IPR> s)
     s.e = 31; s.f[0] = 40; s.f[1] = 39; s.g[1] = 38; s.g[2] = 37;
     s.h[2] = 36; s.h[3] = 35; s.h[4] = 34;
   }
-  #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
   if (!omp_target_is_present (&s.a, d)
       || !omp_target_is_present (s.b, d)
       || !omp_target_is_present (&s.c[1], d)
@@ -135,7 +135,7 @@ foo (S<C, I, L, UCR, CAR, SH, IPR> s)
       || !omp_target_is_present (&s.h, d)
       || !omp_target_is_present (&s.h[2], d))
     abort ();
-  #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
   if (sep
       && (omp_target_is_present (&s.a, d)
 	  || omp_target_is_present (s.b, d)
diff --git a/libgomp/testsuite/libgomp.c++/target-21.C b/libgomp/testsuite/libgomp.c++/target-21.C
index 21a2f299bbb..da17b5745de 100644
--- a/libgomp/testsuite/libgomp.c++/target-21.C
+++ b/libgomp/testsuite/libgomp.c++/target-21.C
@@ -7,7 +7,7 @@ void
 foo (S s)
 {
   int err;
-  #pragma omp target map (s.x[0:N], s.y[0:N]) map (s.t.t[16:3]) map (from: err)
+  #pragma omp target map (s.x[0:N], s.y, s.y[0:N]) map (s.t.t[16:3]) map (from: err)
   {
     err = s.x[2] != 28 || s.y[2] != 37 || s.t.t[17] != 81;
     s.x[2]++;
@@ -38,7 +38,7 @@ void
 foo2 (S &s)
 {
   int err;
-  #pragma omp target map (s.x[N:10], s.y[N:10]) map (from: err) map (s.t.t[N+16:N+3])
+  #pragma omp target map (s.x[N:10], s.y, s.y[N:10]) map (from: err) map (s.t.t[N+16:N+3])
   {
     err = s.x[2] != 30 || s.y[2] != 38 || s.t.t[17] != 81;
     s.x[2]++;
@@ -69,7 +69,7 @@ void
 foo3 (U s)
 {
   int err;
-  #pragma omp target map (s.x[0:10], s.y[0:10]) map (from: err) map (s.t.t[16:3])
+  #pragma omp target map (s.x[0:10], s.y, s.y[0:10]) map (from: err) map (s.t.t[16:3])
   {
     err = s.x[2] != 32 || s.y[2] != 39 || s.t.t[17] != 82;
     s.x[2]++;
@@ -100,7 +100,7 @@ void
 foo4 (U &s)
 {
   int err;
-  #pragma omp target map (s.x[0:10], s.y[0:10]) map (from: err) map (s.t.t[16:3])
+  #pragma omp target map (s.x[0:10], s.y, s.y[0:10]) map (from: err) map (s.t.t[16:3])
   {
     err = s.x[2] != 34 || s.y[2] != 40 || s.t.t[17] != 82;
     s.x[2]++;
diff --git a/libgomp/testsuite/libgomp.c++/target-23.C b/libgomp/testsuite/libgomp.c++/target-23.C
index d4f9ff3e983..63d343624b0 100644
--- a/libgomp/testsuite/libgomp.c++/target-23.C
+++ b/libgomp/testsuite/libgomp.c++/target-23.C
@@ -16,13 +16,13 @@ main (void)
     s->data[i] = 0;
 
   #pragma omp target enter data map(to: s)
-  #pragma omp target enter data map(to: s->data[:SZ])
+  #pragma omp target enter data map(to: s->data, s->data[:SZ])
   #pragma omp target
   {
     for (int i = 0; i < SZ; i++)
       s->data[i] = i;
   }
-  #pragma omp target exit data map(from: s->data[:SZ])
+  #pragma omp target exit data map(from: s->data, s->data[:SZ])
   #pragma omp target exit data map(from: s)
 
   for (int i = 0; i < SZ; i++)
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-2.c b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-2.c
new file mode 100644
index 00000000000..974a9786c3f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-2.c
@@ -0,0 +1,46 @@
+#include <stdlib.h>
+
+#define N 10
+
+struct S
+{
+  int a, b;
+  int *ptr;
+  int c, d;
+};
+
+int
+main (void)
+{
+  struct S a;
+  a.ptr = (int *) malloc (sizeof (int) * N);
+
+  for (int i = 0; i < N; i++)
+    a.ptr[i] = 0;
+
+  #pragma omp target enter data map(to: a.ptr, a.ptr[:N])
+
+  #pragma omp target
+  for (int i = 0; i < N; i++)
+    a.ptr[i] += 1;
+
+  #pragma omp target update from(a.ptr[:N])
+
+  for (int i = 0; i < N; i++)
+    if (a.ptr[i] != 1)
+      abort ();
+
+  #pragma omp target map(a.ptr[:N])
+  for (int i = 0; i < N; i++)
+    a.ptr[i] += 1;
+
+  #pragma omp target update from(a.ptr[:N])
+
+  for (int i = 0; i < N; i++)
+    if (a.ptr[i] != 2)
+      abort ();
+
+  #pragma omp target exit data map(from:a.ptr, a.ptr[:N])
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/target-23.c b/libgomp/testsuite/libgomp.c/target-23.c
index fb1532a07b2..d56b13acf82 100644
--- a/libgomp/testsuite/libgomp.c/target-23.c
+++ b/libgomp/testsuite/libgomp.c/target-23.c
@@ -8,7 +8,7 @@ main ()
   int u[10] = { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9 }, err = 0;
   struct S s = { 9, u + 3, { 10, 11, 12, 13, 14 } };
   int *v = u + 4;
-  #pragma omp target enter data map (to: s.s, s.u[0:5]) map (alloc: s.v[1:3])
+  #pragma omp target enter data map (to: s.s, s.u, s.u[0:5]) map (alloc: s.v[1:3])
   s.s++;
   u[3]++;
   s.v[1]++;
diff --git a/libgomp/testsuite/libgomp.c/target-29.c b/libgomp/testsuite/libgomp.c/target-29.c
index e5095a1b6b8..4a286649811 100644
--- a/libgomp/testsuite/libgomp.c/target-29.c
+++ b/libgomp/testsuite/libgomp.c/target-29.c
@@ -14,7 +14,7 @@ foo (struct S s)
     d = id;
 
   int err;
-  #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d[-2:3]) map(to: sep) map(from: err)
+  #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d, s.d[-2:3]) map(to: sep) map(from: err)
   {
     err = s.a != 11 || s.b[0] != 12 || s.b[1] != 13;
     err |= s.c[1] != 15 || s.c[2] != 16 || s.d[-2] != 18 || s.d[-1] != 19 || s.d[0] != 20;
@@ -35,7 +35,7 @@ foo (struct S s)
 	  || omp_target_is_present (s.d, d)
 	  || omp_target_is_present (&s.d[-2], d)))
     abort ();
-  #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3])
+  #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
   {
     if (!omp_target_is_present (&s.a, d)
 	|| !omp_target_is_present (s.b, d)
@@ -43,15 +43,15 @@ foo (struct S s)
 	|| !omp_target_is_present (s.d, d)
 	|| !omp_target_is_present (&s.d[-2], d))
       abort ();
-    #pragma omp target update to(s.a, s.b, s.c[1:2], s.d[-2:3])
-    #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3]) map(from: err)
+    #pragma omp target update to(s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
+    #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3]) map(from: err)
     {
       err = s.a != 50 || s.b[0] != 49 || s.b[1] != 48;
       err |= s.c[1] != 47 || s.c[2] != 46 || s.d[-2] != 45 || s.d[-1] != 44 || s.d[0] != 43;
       s.a = 17; s.b[0] = 18; s.b[1] = 19;
       s.c[1] = 20; s.c[2] = 21; s.d[-2] = 22; s.d[-1] = 23; s.d[0] = 24;
     }
-    #pragma omp target update from(s.a, s.b, s.c[1:2], s.d[-2:3])
+    #pragma omp target update from(s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
   }
   if (sep
       && (omp_target_is_present (&s.a, d)
@@ -66,29 +66,29 @@ foo (struct S s)
   if (err) abort ();
   s.a = 33; s.b[0] = 34; s.b[1] = 35;
   s.c[1] = 36; s.c[2] = 37; s.d[-2] = 38; s.d[-1] = 39; s.d[0] = 40;
-  #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3])
+  #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
   if (!omp_target_is_present (&s.a, d)
       || !omp_target_is_present (s.b, d)
       || !omp_target_is_present (&s.c[1], d)
       || !omp_target_is_present (s.d, d)
       || !omp_target_is_present (&s.d[-2], d))
     abort ();
-  #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d[-2:3])
-  #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3]) map(from: err)
+  #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
+  #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3]) map(from: err)
   {
     err = s.a != 33 || s.b[0] != 34 || s.b[1] != 35;
     err |= s.c[1] != 36 || s.c[2] != 37 || s.d[-2] != 38 || s.d[-1] != 39 || s.d[0] != 40;
     s.a = 49; s.b[0] = 48; s.b[1] = 47;
     s.c[1] = 46; s.c[2] = 45; s.d[-2] = 44; s.d[-1] = 43; s.d[0] = 42;
   }
-  #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d[-2:3])
+  #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
   if (!omp_target_is_present (&s.a, d)
       || !omp_target_is_present (s.b, d)
       || !omp_target_is_present (&s.c[1], d)
       || !omp_target_is_present (s.d, d)
       || !omp_target_is_present (&s.d[-2], d))
     abort ();
-  #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d[-2:3])
+  #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
   if (sep
       && (omp_target_is_present (&s.a, d)
 	  || omp_target_is_present (s.b, d)
