diff --git a/gcc/c-family/c-common.h b/gcc/c-family/c-common.h
index 01ec9a739458..649c4bff7f2e 100644
--- a/gcc/c-family/c-common.h
+++ b/gcc/c-family/c-common.h
@@ -1349,7 +1349,7 @@ public:
   bool maybe_zero_length_array_section (tree);
 
   tree * expand_array_base (tree *, vec<omp_addr_token *> &, tree, unsigned *,
-			    c_omp_region_type, bool);
+			    c_omp_region_type);
   tree * expand_component_selector (tree *, vec<omp_addr_token *> &, tree,
 				    unsigned *, c_omp_region_type);
   tree * expand_map_clause (tree *, tree, vec<omp_addr_token *> &,
diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc
index 0de3d350d023..015a5882e36a 100644
--- a/gcc/c-family/c-omp.cc
+++ b/gcc/c-family/c-omp.cc
@@ -3244,7 +3244,9 @@ c_omp_address_inspector::map_supported_p ()
 	 || TREE_CODE (t) == POINTER_PLUS_EXPR
 	 || TREE_CODE (t) == NON_LVALUE_EXPR
 	 || TREE_CODE (t) == OMP_ARRAY_SECTION
-	 || TREE_CODE (t) == NOP_EXPR)
+	 || TREE_CODE (t) == NOP_EXPR
+	 || TREE_CODE (t) == VIEW_CONVERT_EXPR
+	 || TREE_CODE (t) == ADDR_EXPR)
     if (TREE_CODE (t) == COMPOUND_EXPR)
       t = TREE_OPERAND (t, 1);
     else
@@ -3384,21 +3386,95 @@ omp_expand_access_chain (tree *pc, tree expr,
   return pc;
 }
 
+static tree *
+omp_expand_grid_dim (location_t loc, tree *pc, tree decl)
+{
+  if (TREE_CODE (decl) == OMP_ARRAY_SECTION)
+    pc = omp_expand_grid_dim (loc, pc, TREE_OPERAND (decl, 0));
+  else
+    return pc;
+
+  tree c = *pc;
+  tree low_bound = TREE_OPERAND (decl, 1);
+  tree length = TREE_OPERAND (decl, 2);
+  tree stride = TREE_OPERAND (decl, 3);
+
+  tree cd = build_omp_clause (loc, OMP_CLAUSE_MAP);
+  OMP_CLAUSE_SET_MAP_KIND (cd, GOMP_MAP_GRID_DIM);
+  OMP_CLAUSE_DECL (cd) = unshare_expr (low_bound);
+  OMP_CLAUSE_SIZE (cd) = unshare_expr (length);
+
+  if (stride && !integer_onep (stride))
+    {
+      tree cs = build_omp_clause (loc, OMP_CLAUSE_MAP);
+      OMP_CLAUSE_SET_MAP_KIND (cs, GOMP_MAP_GRID_STRIDE);
+      OMP_CLAUSE_DECL (cs) = unshare_expr (stride);
+
+      OMP_CLAUSE_CHAIN (cs) = OMP_CLAUSE_CHAIN (c);
+      OMP_CLAUSE_CHAIN (cd) = cs;
+      OMP_CLAUSE_CHAIN (c) = cd;
+      pc = &OMP_CLAUSE_CHAIN (cd);
+    }
+  else
+    {
+      OMP_CLAUSE_CHAIN (cd) = OMP_CLAUSE_CHAIN (c);
+      OMP_CLAUSE_CHAIN (c) = cd;
+      pc = &OMP_CLAUSE_CHAIN (c);
+    }
+
+  return pc;
+}
+
+tree *
+omp_handle_noncontig_array (location_t loc, tree *pc, tree c, tree base)
+{
+  tree type;
+
+  if (POINTER_TYPE_P (TREE_TYPE (base)))
+    type = TREE_TYPE (TREE_TYPE (base));
+  else
+    type = strip_array_types (TREE_TYPE (base));
+
+  tree c_map = build_omp_clause (loc, OMP_CLAUSE_MAP);
+
+  OMP_CLAUSE_DECL (c_map) = unshare_expr (base);
+  /* Use the element size (or pointed-to type size) here.  */
+  OMP_CLAUSE_SIZE (c_map) = TYPE_SIZE_UNIT (type);
+
+  switch (OMP_CLAUSE_CODE (c))
+    {
+    case OMP_CLAUSE_TO:
+      OMP_CLAUSE_SET_MAP_KIND (c_map, GOMP_MAP_TO_GRID);
+      break;
+    case OMP_CLAUSE_FROM:
+      OMP_CLAUSE_SET_MAP_KIND (c_map, GOMP_MAP_FROM_GRID);
+      break;
+    default:
+      gcc_unreachable ();
+    }
+
+  OMP_CLAUSE_CHAIN (c_map) = OMP_CLAUSE_CHAIN (c);
+
+  *pc = c_map;
+
+  return omp_expand_grid_dim (loc, pc, OMP_CLAUSE_DECL (c));
+}
+
 /* Translate "array_base_decl access_method" to OMP mapping clauses.  */
 
 tree *
 c_omp_address_inspector::expand_array_base (tree *pc,
 					    vec<omp_addr_token *> &addr_tokens,
 					    tree expr, unsigned *idx,
-					    c_omp_region_type ort,
-					    bool decl_p)
+					    c_omp_region_type ort)
 {
   using namespace omp_addr_tokenizer;
   tree c = *pc;
   location_t loc = OMP_CLAUSE_LOCATION (c);
   int i = *idx;
   tree decl = addr_tokens[i + 1]->expr;
-  bool declare_target_p = (decl_p
+  bool decl_p = DECL_P (decl);
+  bool declare_target_p = (DECL_P (decl)
 			   && is_global_var (decl)
 			   && lookup_attribute ("omp declare target",
 						DECL_ATTRIBUTES (decl)));
@@ -3410,6 +3486,7 @@ c_omp_address_inspector::expand_array_base (tree *pc,
   unsigned consume_tokens = 2;
   bool target = (ort & C_ORT_TARGET) != 0;
   bool openmp = (ort & C_ORT_OMP) != 0;
+  unsigned acc = i + 1;
 
   gcc_assert (i == 0);
 
@@ -3422,7 +3499,15 @@ c_omp_address_inspector::expand_array_base (tree *pc,
       return pc;
     }
 
-  switch (addr_tokens[i + 1]->u.access_kind)
+  if (!map_p && chain_p)
+    {
+      /* See comment in c_omp_address_inspector::expand_component_selector.  */
+      while (acc + 1 < addr_tokens.length ()
+	     && addr_tokens[acc + 1]->type == ACCESS_METHOD)
+	acc++;
+    }
+
+  switch (addr_tokens[acc]->u.access_kind)
     {
     case ACCESS_DIRECT:
       if (decl_p && !target)
@@ -3663,6 +3748,40 @@ c_omp_address_inspector::expand_array_base (tree *pc,
       }
       break;
 
+    case ACCESS_NONCONTIG_ARRAY:
+      {
+	gcc_assert (!map_p);
+
+	tree base = addr_tokens[acc]->expr;
+
+	if (decl_p)
+	  c_common_mark_addressable_vec (base);
+
+	pc = omp_handle_noncontig_array (loc, pc, c, base);
+	consume_tokens = (acc + 1) - i;
+	chain_p = false;
+      }
+      break;
+
+    case ACCESS_NONCONTIG_REF_TO_ARRAY:
+      {
+	gcc_assert (!map_p);
+
+	if (decl_p)
+	  c_common_mark_addressable_vec (addr_tokens[acc]->expr);
+
+	/* Or here.  */
+	gcc_assert (!chain_p);
+
+	tree base = addr_tokens[i + 1]->expr;
+	base = convert_from_reference (base);
+
+	pc = omp_handle_noncontig_array (loc, pc, c, base);
+	consume_tokens = (acc + 1) - i;
+	chain_p = false;
+      }
+      break;
+
     default:
       *idx = i + consume_tokens;
       return NULL;
@@ -3719,8 +3838,27 @@ c_omp_address_inspector::expand_component_selector (tree *pc,
   bool chain_p = omp_access_chain_p (addr_tokens, i + 1);
   bool map_p = OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP;
   bool target = (ort & C_ORT_TARGET) != 0;
+  unsigned acc = i + 1;
 
-  switch (addr_tokens[i + 1]->u.access_kind)
+  if (!map_p && chain_p)
+    {
+      /* We have a non-map clause (i.e. to/from for an "update" directive),
+	 and we might have a noncontiguous array section at the end of a
+	 chain of other accesses, e.g. pointer indirections like this:
+
+	   struct_base_decl access_pointer access_pointer component_selector
+	     access_pointer access_pointer access_noncontig_array
+
+	 We only need to process the last access in this case, so skip
+	 over previous accesses.  */
+
+      while (acc + 1 < addr_tokens.length ()
+	     && addr_tokens[acc + 1]->type == ACCESS_METHOD)
+	acc++;
+      chain_p = false;
+    }
+
+  switch (addr_tokens[acc]->u.access_kind)
     {
     case ACCESS_DIRECT:
     case ACCESS_INDEXED_ARRAY:
@@ -3730,7 +3868,7 @@ c_omp_address_inspector::expand_component_selector (tree *pc,
       {
 	/* Copy the referenced object.  Note that we also do this for !MAP_P
 	   clauses.  */
-	tree obj = convert_from_reference (addr_tokens[i + 1]->expr);
+	tree obj = convert_from_reference (addr_tokens[acc]->expr);
 	OMP_CLAUSE_DECL (c) = obj;
 	OMP_CLAUSE_SIZE (c) = TYPE_SIZE_UNIT (TREE_TYPE (obj));
 
@@ -3739,7 +3877,7 @@ c_omp_address_inspector::expand_component_selector (tree *pc,
 
 	c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
 	OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
-	OMP_CLAUSE_DECL (c2) = addr_tokens[i + 1]->expr;
+	OMP_CLAUSE_DECL (c2) = addr_tokens[acc]->expr;
 	OMP_CLAUSE_SIZE (c2) = size_zero_node;
       }
       break;
@@ -3750,15 +3888,15 @@ c_omp_address_inspector::expand_component_selector (tree *pc,
 	  break;
 
 	tree virtual_origin
-	  = convert_from_reference (addr_tokens[i + 1]->expr);
+	  = convert_from_reference (addr_tokens[acc]->expr);
 	virtual_origin = build_fold_addr_expr (virtual_origin);
 	virtual_origin = fold_convert_loc (loc, ptrdiff_type_node,
 					   virtual_origin);
-	tree data_addr = omp_accessed_addr (addr_tokens, i + 1, expr);
+	tree data_addr = omp_accessed_addr (addr_tokens, acc, expr);
 
 	c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
 	OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
-	OMP_CLAUSE_DECL (c2) = addr_tokens[i + 1]->expr;
+	OMP_CLAUSE_DECL (c2) = addr_tokens[acc]->expr;
 	OMP_CLAUSE_SIZE (c2)
 	  = fold_build2_loc (loc, MINUS_EXPR, ptrdiff_type_node,
 			     fold_convert_loc (loc, ptrdiff_type_node,
@@ -3775,12 +3913,12 @@ c_omp_address_inspector::expand_component_selector (tree *pc,
 
 	tree virtual_origin
 	  = fold_convert_loc (loc, ptrdiff_type_node,
-			      addr_tokens[i + 1]->expr);
-	tree data_addr = omp_accessed_addr (addr_tokens, i + 1, expr);
+			      addr_tokens[acc]->expr);
+	tree data_addr = omp_accessed_addr (addr_tokens, acc, expr);
 
 	c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
 	OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
-	OMP_CLAUSE_DECL (c2) = addr_tokens[i + 1]->expr;
+	OMP_CLAUSE_DECL (c2) = addr_tokens[acc]->expr;
 	OMP_CLAUSE_SIZE (c2)
 	  = fold_build2_loc (loc, MINUS_EXPR, ptrdiff_type_node,
 			     fold_convert_loc (loc, ptrdiff_type_node,
@@ -3795,10 +3933,10 @@ c_omp_address_inspector::expand_component_selector (tree *pc,
 	if (!map_p)
 	  break;
 
-	tree ptr = convert_from_reference (addr_tokens[i + 1]->expr);
+	tree ptr = convert_from_reference (addr_tokens[acc]->expr);
 	tree virtual_origin = fold_convert_loc (loc, ptrdiff_type_node,
 						ptr);
-	tree data_addr = omp_accessed_addr (addr_tokens, i + 1, expr);
+	tree data_addr = omp_accessed_addr (addr_tokens, acc, expr);
 
 	/* Attach the pointer...  */
 	c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
@@ -3813,13 +3951,38 @@ c_omp_address_inspector::expand_component_selector (tree *pc,
 	/* ...and also the reference.  */
 	c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
 	OMP_CLAUSE_SET_MAP_KIND (c3, GOMP_MAP_ATTACH_DETACH);
-	OMP_CLAUSE_DECL (c3) = addr_tokens[i + 1]->expr;
+	OMP_CLAUSE_DECL (c3) = addr_tokens[acc]->expr;
 	OMP_CLAUSE_SIZE (c3) = size_zero_node;
       }
       break;
 
+    case ACCESS_NONCONTIG_ARRAY:
+      {
+	gcc_assert (!map_p);
+
+	/* We don't expect to see further accesses here.  */
+	gcc_assert (!chain_p);
+
+	pc = omp_handle_noncontig_array (loc, pc, c, addr_tokens[acc]->expr);
+      }
+      break;
+
+    case ACCESS_NONCONTIG_REF_TO_ARRAY:
+      {
+	gcc_assert (!map_p);
+
+	/* Or here.  */
+	gcc_assert (!chain_p);
+
+	tree base = addr_tokens[acc]->expr;
+	base = convert_from_reference (base);
+
+	pc = omp_handle_noncontig_array (loc, pc, c, base);
+      }
+      break;
+
     default:
-      *idx = i + 2;
+      *idx = acc + 1;
       return NULL;
     }
 
@@ -3837,8 +4000,7 @@ c_omp_address_inspector::expand_component_selector (tree *pc,
       pc = &OMP_CLAUSE_CHAIN (c);
     }
 
-  i += 2;
-  *idx = i;
+  *idx = acc + 1;
 
   if (target && chain_p && map_p)
     return omp_expand_access_chain (pc, expr, addr_tokens, idx);
@@ -3870,7 +4032,7 @@ c_omp_address_inspector::expand_map_clause (tree *pc, tree expr,
 	  && addr_tokens[i]->u.structure_base_kind == BASE_DECL
 	  && addr_tokens[i + 1]->type == ACCESS_METHOD)
 	{
-	  pc = expand_array_base (pc, addr_tokens, expr, &i, ort, true);
+	  pc = expand_array_base (pc, addr_tokens, expr, &i, ort);
 	  if (pc == NULL)
 	    return NULL;
 	}
@@ -3879,7 +4041,7 @@ c_omp_address_inspector::expand_map_clause (tree *pc, tree expr,
 	       && addr_tokens[i]->u.structure_base_kind == BASE_ARBITRARY_EXPR
 	       && addr_tokens[i + 1]->type == ACCESS_METHOD)
 	{
-	  pc = expand_array_base (pc, addr_tokens, expr, &i, ort, false);
+	  pc = expand_array_base (pc, addr_tokens, expr, &i, ort);
 	  if (pc == NULL)
 	    return NULL;
 	}
diff --git a/gcc/c-family/c-pretty-print.cc b/gcc/c-family/c-pretty-print.cc
index d9954bd2b951..bfd5e2e14354 100644
--- a/gcc/c-family/c-pretty-print.cc
+++ b/gcc/c-family/c-pretty-print.cc
@@ -1623,6 +1623,11 @@ c_pretty_printer::postfix_expression (tree e)
       pp_colon (this);
       if (TREE_OPERAND (e, 2))
 	expression (TREE_OPERAND (e, 2));
+      if (TREE_OPERAND (e, 3))
+	{
+	  pp_colon (this);
+	  expression (TREE_OPERAND (e, 3));
+	}
       pp_c_right_bracket (this);
       break;
 
diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 0d5013a3f06a..bdf5ee9fc30d 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -11224,7 +11224,7 @@ c_parser_postfix_expression_after_primary (c_parser *parser,
 	      start = expr.get_start ();
 	      finish = parser->tokens_buf[0].location;
 	      expr.value = build_omp_array_section (op_loc, expr.value, idx,
-						    len);
+						    len, NULL_TREE /* fixme */);
 	      set_c_expr_source_range (&expr, start, finish);
 	      expr.original_code = ERROR_MARK;
 	      expr.original_type = NULL;
@@ -13760,11 +13760,11 @@ c_parser_oacc_wait_list (c_parser *parser, location_t clause_loc, tree list)
 
 struct omp_dim
 {
-  tree low_bound, length;
+  tree low_bound, length, stride;
   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) {}
+  omp_dim (tree lb, tree len, tree str, location_t lo, bool nc)
+  : low_bound (lb), length (len), stride (str), loc (lo), no_colon (nc) {}
 };
 
 static tree
@@ -13892,7 +13892,9 @@ c_parser_omp_variable_list (c_parser *parser,
 		{
 		  tree low_bound = TREE_OPERAND (decl, 1);
 		  tree length = TREE_OPERAND (decl, 2);
-		  dims.safe_push (omp_dim (low_bound, length, loc, false));
+		  tree stride = TREE_OPERAND (decl, 3);
+		  dims.safe_push (omp_dim (low_bound, length, stride, loc,
+					   false));
 		  decl = TREE_OPERAND (decl, 0);
 		}
 
@@ -13908,21 +13910,22 @@ c_parser_omp_variable_list (c_parser *parser,
 		  else if (TREE_CODE (decl) == INDIRECT_REF)
 		    {
 		      dims.safe_push (omp_dim (integer_zero_node,
-					       integer_one_node, loc, true));
+					       integer_one_node, NULL_TREE, 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));
+		      dims.safe_push (omp_dim (index, integer_one_node,
+					       NULL_TREE, loc, true));
 		      decl = TREE_OPERAND (decl, 0);
 		    }
 		}
 
 	      for (int i = dims.length () - 1; i >= 0; i--)
 		decl = build_omp_array_section (loc,  decl, dims[i].low_bound,
-						dims[i].length);
+						dims[i].length, dims[i].stride);
 	    }
 	  else if (TREE_CODE (decl) == INDIRECT_REF)
 	    {
@@ -13932,7 +13935,7 @@ c_parser_omp_variable_list (c_parser *parser,
 	      STRIP_NOPS (decl);
 
 	      decl = build_omp_array_section (loc, decl, integer_zero_node,
-					      integer_one_node);
+					      integer_one_node, NULL_TREE);
 	    }
 	  else if (TREE_CODE (decl) == ARRAY_REF)
 	    {
@@ -13941,7 +13944,8 @@ c_parser_omp_variable_list (c_parser *parser,
 	      decl = TREE_OPERAND (decl, 0);
 	      STRIP_NOPS (decl);
 
-	      decl = build_omp_array_section (loc, decl, idx, integer_one_node);
+	      decl = build_omp_array_section (loc, decl, idx, integer_one_node,
+					      NULL_TREE);
 	    }
 	  else if (TREE_CODE (decl) == NON_LVALUE_EXPR
 		   || CONVERT_EXPR_P (decl))
@@ -14095,7 +14099,8 @@ c_parser_omp_variable_list (c_parser *parser,
 		      break;
 		    }
 
-		  dims.safe_push (omp_dim (low_bound, length, loc, no_colon));
+		  dims.safe_push (omp_dim (low_bound, length, NULL_TREE, loc,
+				  no_colon));
 		}
 
 	      if (t != error_mark_node)
@@ -14119,7 +14124,8 @@ c_parser_omp_variable_list (c_parser *parser,
 		    for (unsigned i = 0; i < dims.length (); i++)
 		      t = build_omp_array_section (clause_loc, t,
 						   dims[i].low_bound,
-						   dims[i].length);
+						   dims[i].length,
+						   dims[i].stride);
 		}
 
 	      if ((kind == OMP_CLAUSE_DEPEND || kind == OMP_CLAUSE_AFFINITY)
diff --git a/gcc/c/c-tree.h b/gcc/c/c-tree.h
index 964623feb472..4ff5c12ff7b6 100644
--- a/gcc/c/c-tree.h
+++ b/gcc/c/c-tree.h
@@ -752,7 +752,7 @@ extern tree composite_type (tree, tree);
 extern tree build_component_ref (location_t, tree, tree, location_t,
 				 location_t);
 extern tree build_array_ref (location_t, tree, tree);
-extern tree build_omp_array_section (location_t, tree, tree, tree);
+extern tree build_omp_array_section (location_t, tree, tree, tree, tree);
 extern tree build_external_ref (location_t, tree, bool, tree *);
 extern void pop_maybe_used (bool);
 extern struct c_expr c_expr_sizeof_expr (location_t, struct c_expr);
diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc
index 12be66bd49be..2c7c802175a5 100644
--- a/gcc/c/c-typeck.cc
+++ b/gcc/c/c-typeck.cc
@@ -2046,6 +2046,8 @@ mark_exp_read (tree exp)
 	mark_exp_read (TREE_OPERAND (exp, 1));
       if (TREE_OPERAND (exp, 2))
 	mark_exp_read (TREE_OPERAND (exp, 2));
+      if (TREE_OPERAND (exp, 3))
+	mark_exp_read (TREE_OPERAND (exp, 3));
       break;
     default:
       break;
@@ -2931,7 +2933,8 @@ build_array_ref (location_t loc, tree array, tree index)
    instead.  */
 
 tree
-build_omp_array_section (location_t loc, tree array, tree index, tree length)
+build_omp_array_section (location_t loc, tree array, tree index, tree length,
+			 tree stride)
 {
   tree idxtype;
 
@@ -2968,7 +2971,8 @@ build_omp_array_section (location_t loc, tree array, tree index, tree length)
   else
     sectype = build_array_type (eltype, idxtype);
 
-  return build3_loc (loc, OMP_ARRAY_SECTION, sectype, array, index, length);
+  return build4_loc (loc, OMP_ARRAY_SECTION, sectype, array, index, length,
+		     stride);
 }
 
 
@@ -13694,7 +13698,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 			     bool &maybe_zero_len, unsigned int &first_non_one,
 			     enum c_omp_region_type ort)
 {
-  tree ret, low_bound, length, type;
+  tree ret, low_bound, length, stride, type;
   bool openacc = (ort & C_ORT_ACC) != 0;
   if (TREE_CODE (t) != OMP_ARRAY_SECTION)
     {
@@ -13779,8 +13783,11 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
   type = TREE_TYPE (ret);
   low_bound = TREE_OPERAND (t, 1);
   length = TREE_OPERAND (t, 2);
+  stride = TREE_OPERAND (t, 3);
 
-  if (low_bound == error_mark_node || length == error_mark_node)
+  if (low_bound == error_mark_node
+      || length == error_mark_node
+      || stride == error_mark_node)
     return error_mark_node;
 
   if (low_bound && !INTEGRAL_TYPE_P (TREE_TYPE (low_bound)))
@@ -13797,6 +13804,13 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 		length);
       return error_mark_node;
     }
+  if (stride && !INTEGRAL_TYPE_P (TREE_TYPE (stride)))
+    {
+      error_at (OMP_CLAUSE_LOCATION (c),
+		"stride %qE of array section does not have integral type",
+		stride);
+      return error_mark_node;
+    }
   if (low_bound
       && TREE_CODE (low_bound) == INTEGER_CST
       && TYPE_PRECISION (TREE_TYPE (low_bound))
@@ -14011,7 +14025,9 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 	       d = TREE_OPERAND (d, 0))
 	    {
 	      tree d_length = TREE_OPERAND (d, 2);
-	      if (d_length == NULL_TREE || !integer_onep (d_length))
+	      tree d_stride = TREE_OPERAND (d, 3);
+	      if (d_length == NULL_TREE || !integer_onep (d_length)
+		  || (d_stride && !integer_onep (d_stride)))
 		{
 		  error_at (OMP_CLAUSE_LOCATION (c),
 			    "array section is not contiguous in %qs clause",
diff --git a/gcc/cp/cp-objcp-common.cc b/gcc/cp/cp-objcp-common.cc
index 6ed382662811..36d4de3a8e16 100644
--- a/gcc/cp/cp-objcp-common.cc
+++ b/gcc/cp/cp-objcp-common.cc
@@ -530,6 +530,7 @@ cp_common_init_ts (void)
   MARK_TS_EXP (OFFSET_REF);
   MARK_TS_EXP (PSEUDO_DTOR_EXPR);
   MARK_TS_EXP (REINTERPRET_CAST_EXPR);
+  MARK_TS_EXP (OMP_ARRAYSHAPE_CAST_EXPR);
   MARK_TS_EXP (SCOPE_REF);
   MARK_TS_EXP (STATIC_CAST_EXPR);
   MARK_TS_EXP (STMT_EXPR);
diff --git a/gcc/cp/cp-tree.def b/gcc/cp/cp-tree.def
index d8aab10549f1..940377f51e2a 100644
--- a/gcc/cp/cp-tree.def
+++ b/gcc/cp/cp-tree.def
@@ -256,6 +256,7 @@ DEFTREECODE (REINTERPRET_CAST_EXPR, "reinterpret_cast_expr", tcc_unary, 1)
 DEFTREECODE (CONST_CAST_EXPR, "const_cast_expr", tcc_unary, 1)
 DEFTREECODE (STATIC_CAST_EXPR, "static_cast_expr", tcc_unary, 1)
 DEFTREECODE (DYNAMIC_CAST_EXPR, "dynamic_cast_expr", tcc_unary, 1)
+DEFTREECODE (OMP_ARRAYSHAPE_CAST_EXPR, "omp_arrayshape_cast_expr", tcc_unary, 1)
 DEFTREECODE (IMPLICIT_CONV_EXPR, "implicit_conv_expr", tcc_unary, 1)
 DEFTREECODE (DOTSTAR_EXPR, "dotstar_expr", tcc_expression, 2)
 DEFTREECODE (TYPEID_EXPR, "typeid_expr", tcc_expression, 1)
diff --git a/gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h
index d80b698b2ac8..3fb9c1658f44 100644
--- a/gcc/cp/cp-tree.h
+++ b/gcc/cp/cp-tree.h
@@ -515,6 +515,7 @@ extern GTY(()) tree cp_global_trees[CPTI_MAX];
       OVL_LOOKUP_P (in OVERLOAD)
       LOOKUP_FOUND_P (in RECORD_TYPE, UNION_TYPE, ENUMERAL_TYPE, NAMESPACE_DECL)
       FNDECL_MANIFESTLY_CONST_EVALUATED (in FUNCTION_DECL)
+      DECLTYPE_FOR_OMP_ARRAYSHAPE_CAST (in DECLTYPE_TYPE)
    5: IDENTIFIER_VIRTUAL_P (in IDENTIFIER_NODE)
       FUNCTION_RVALUE_QUALIFIED (in FUNCTION_TYPE, METHOD_TYPE)
       CALL_EXPR_REVERSE_ARGS (in CALL_EXPR, AGGR_INIT_EXPR)
@@ -4863,6 +4864,8 @@ get_vec_init_expr (tree t)
   TREE_LANG_FLAG_2 (DECLTYPE_TYPE_CHECK (NODE))
 #define DECLTYPE_FOR_REF_CAPTURE(NODE) \
   TREE_LANG_FLAG_3 (DECLTYPE_TYPE_CHECK (NODE))
+#define DECLTYPE_FOR_OMP_ARRAYSHAPE_CAST(NODE) \
+  TREE_LANG_FLAG_4 (DECLTYPE_TYPE_CHECK (NODE))
 
 /* Nonzero for VAR_DECL and FUNCTION_DECL node means that `extern' was
    specified in its declaration.  This can also be set for an
@@ -6931,6 +6934,8 @@ extern tree cxx_comdat_group			(tree);
 extern bool cp_missing_noreturn_ok_p		(tree);
 extern bool is_direct_enum_init			(tree, tree);
 extern void initialize_artificial_var		(tree, vec<constructor_elt, va_gc> *);
+extern tree cp_omp_create_arrayshape_type	(location_t, tree,
+						 vec<cp_expr> *);
 extern tree check_var_type			(tree, tree, location_t);
 extern tree reshape_init                        (tree, tree, tsubst_flags_t);
 extern tree next_aggregate_field		(tree);
@@ -6964,7 +6969,8 @@ extern void grokclassfn				(tree, tree,
 						 enum overload_flags);
 extern tree grok_array_decl			(location_t, tree, tree,
 						 vec<tree, va_gc> **, tsubst_flags_t);
-extern tree grok_omp_array_section		(location_t, tree, tree, tree);
+extern tree grok_omp_array_section		(location_t, tree, tree, tree,
+						 tree);
 extern tree delete_sanity			(location_t, tree, tree, bool,
 						 int, tsubst_flags_t);
 extern tree check_classfn			(tree, tree, tree);
@@ -7812,6 +7818,8 @@ extern tree cp_build_vec_convert		(tree, location_t, tree,
 						 tsubst_flags_t);
 extern tree cp_build_bit_cast			(location_t, tree, tree,
 						 tsubst_flags_t);
+extern tree cp_build_omp_arrayshape_cast	(location_t, tree, tree,
+						 tsubst_flags_t);
 extern void start_lambda_scope			(tree decl);
 extern void finish_lambda_scope			(void);
 extern void record_lambda_scope			(tree lambda);
@@ -8063,7 +8071,8 @@ inline tree build_x_binary_op (const op_location_t &loc,
 }
 extern tree build_x_array_ref			(location_t, tree, tree,
 						 tsubst_flags_t);
-extern tree build_omp_array_section		(location_t, tree, tree, tree);
+extern tree build_omp_array_section		(location_t, tree, tree, tree,
+						 tree);
 extern tree build_x_unary_op			(location_t,
 						 enum tree_code, cp_expr,
 						 tree, tsubst_flags_t);
diff --git a/gcc/cp/decl.cc b/gcc/cp/decl.cc
index 72e7d5f9de04..88484ccfe9c8 100644
--- a/gcc/cp/decl.cc
+++ b/gcc/cp/decl.cc
@@ -11552,6 +11552,81 @@ create_array_type_for_decl (tree name, tree type, tree size, location_t loc)
   return build_cplus_array_type (type, itype);
 }
 
+/* Build an anonymous array of SIZE elements of ELTYPE.  */
+
+static tree
+create_anon_array_type (location_t loc, tree eltype, tree size)
+{
+  if (eltype == error_mark_node || size == error_mark_node)
+    return error_mark_node;
+
+  tree itype = compute_array_index_type_loc (loc, NULL_TREE, size,
+					     tf_warning_or_error);
+
+  if (type_uses_auto (eltype)
+      && variably_modified_type_p (itype, /*fn=*/NULL_TREE))
+    {
+      sorry_at (loc, "variable-length array of %<auto%>");
+      return error_mark_node;
+    }
+
+  return build_cplus_array_type (eltype, itype);
+}
+
+/* Derive an array type for an OpenMP array-shaping operator given EXPR, which
+   is an expression that might have array refs or array sections postfixed
+   (e.g. "ptr[0:3:2][3:4]"), and OMP_SHAPE_DIMS, a vector of dimensions.  */
+
+tree
+cp_omp_create_arrayshape_type (location_t loc, tree expr,
+			       vec<cp_expr> *omp_shape_dims)
+{
+  tree type, strip_sections = expr;
+
+  while (TREE_CODE (strip_sections) == OMP_ARRAY_SECTION
+	 || TREE_CODE (strip_sections) == ARRAY_REF)
+    strip_sections = TREE_OPERAND (strip_sections, 0);
+
+  /* Determine the element type, either directly or by using
+     "decltype" of an expression representing an element to
+     figure it out later during template instantiation.  */
+  if (type_dependent_expression_p (expr))
+    {
+      type = cxx_make_type (DECLTYPE_TYPE);
+
+      DECLTYPE_TYPE_EXPR (type)
+	= build_min_nt_loc (loc, INDIRECT_REF, strip_sections);
+      DECLTYPE_FOR_OMP_ARRAYSHAPE_CAST (type) = true;
+      SET_TYPE_STRUCTURAL_EQUALITY (type);
+    }
+  else
+    {
+      type = TREE_TYPE (strip_sections);
+
+      if (TREE_CODE (type) == REFERENCE_TYPE)
+	type = TREE_TYPE (type);
+
+      if (TREE_CODE (type) != POINTER_TYPE)
+	{
+	  error ("OpenMP array shaping operator with non-pointer argument");
+	  return error_mark_node;
+	}
+
+      type = TREE_TYPE (type);
+    }
+
+  int i;
+  cp_expr dim;
+  FOR_EACH_VEC_ELT_REVERSE (*omp_shape_dims, i, dim)
+    {
+      if (!type_dependent_expression_p (dim))
+	dim = fold_convert (sizetype, dim);
+      type = create_anon_array_type (loc, type, dim);
+    }
+
+  return type;
+}
+
 /* Returns the smallest location that is not UNKNOWN_LOCATION.  */
 
 static location_t
diff --git a/gcc/cp/decl2.cc b/gcc/cp/decl2.cc
index fc6fd28046c8..608974aaba82 100644
--- a/gcc/cp/decl2.cc
+++ b/gcc/cp/decl2.cc
@@ -596,43 +596,49 @@ grok_array_decl (location_t loc, tree array_expr, tree index_exp,
 
 tree
 grok_omp_array_section (location_t loc, tree array_expr, tree index,
-			tree length)
+			tree length, tree stride)
 {
   tree orig_array_expr = array_expr;
   tree orig_index = index;
   tree orig_length = length;
+  tree orig_stride = stride;
 
   if (error_operand_p (array_expr)
       || error_operand_p (index)
-      || error_operand_p (length))
+      || error_operand_p (length)
+      || error_operand_p (stride))
     return error_mark_node;
 
   if (processing_template_decl)
     {
       if (type_dependent_expression_p (array_expr)
 	  || type_dependent_expression_p (index)
-	  || type_dependent_expression_p (length))
+	  || type_dependent_expression_p (length)
+	  || type_dependent_expression_p (stride))
 	return build_min_nt_loc (loc, OMP_ARRAY_SECTION, array_expr, index,
-				 length);
+				 length, stride);
       array_expr = build_non_dependent_expr (array_expr);
       if (index)
 	index = build_non_dependent_expr (index);
       if (length)
 	length = build_non_dependent_expr (length);
+      if (stride)
+	stride = build_non_dependent_expr (stride);
     }
 
   index = fold_non_dependent_expr (index);
   length = fold_non_dependent_expr (length);
+  stride = fold_non_dependent_expr (stride);
 
   /* NOTE: We can pass through invalidly-typed index/length fields
      here (e.g. if the user tries to use a floating-point index/length).
      This is diagnosed later in semantics.cc:handle_omp_array_sections_1.  */
 
-  tree expr = build_omp_array_section (loc, array_expr, index, length);
+  tree expr = build_omp_array_section (loc, array_expr, index, length, stride);
 
   if (processing_template_decl)
     expr = build_min_non_dep (OMP_ARRAY_SECTION, expr, orig_array_expr,
-			      orig_index, orig_length);
+			      orig_index, orig_length, orig_stride);
   return expr;
 }
 
@@ -2608,6 +2614,7 @@ min_vis_expr_r (tree *tp, int */*walk_subtrees*/, void *data)
     case REINTERPRET_CAST_EXPR:
     case CONST_CAST_EXPR:
     case DYNAMIC_CAST_EXPR:
+    case OMP_ARRAYSHAPE_CAST_EXPR:
     case NEW_EXPR:
     case CONSTRUCTOR:
     case LAMBDA_EXPR:
diff --git a/gcc/cp/error.cc b/gcc/cp/error.cc
index 52b66975d238..f45ac031c8bb 100644
--- a/gcc/cp/error.cc
+++ b/gcc/cp/error.cc
@@ -2528,6 +2528,11 @@ dump_expr (cxx_pretty_printer *pp, tree t, int flags)
       dump_expr (pp, TREE_OPERAND (t, 1), flags);
       pp_colon (pp);
       dump_expr (pp, TREE_OPERAND (t, 2), flags);
+      if (TREE_OPERAND (t, 3))
+	{
+	  pp_colon (pp);
+	  dump_expr (pp, TREE_OPERAND (t, 3), flags);
+	}
       pp_cxx_right_bracket (pp);
       break;
 
diff --git a/gcc/cp/mangle.cc b/gcc/cp/mangle.cc
index 5789adcf680f..43606da24307 100644
--- a/gcc/cp/mangle.cc
+++ b/gcc/cp/mangle.cc
@@ -3580,6 +3580,7 @@ write_expression (tree expr)
 	case REINTERPRET_CAST_EXPR:
 	case STATIC_CAST_EXPR:
 	case CONST_CAST_EXPR:
+	case OMP_ARRAYSHAPE_CAST_EXPR:
 	  write_type (TREE_TYPE (expr));
 	  write_expression (TREE_OPERAND (expr, 0));
 	  break;
diff --git a/gcc/cp/operators.def b/gcc/cp/operators.def
index 27b8559c8983..92bbe8d8e11a 100644
--- a/gcc/cp/operators.def
+++ b/gcc/cp/operators.def
@@ -134,6 +134,7 @@ DEF_OPERATOR (NULL, DYNAMIC_CAST_EXPR, "dc", OVL_OP_FLAG_UNARY)
 DEF_OPERATOR (NULL, REINTERPRET_CAST_EXPR, "rc", OVL_OP_FLAG_UNARY)
 DEF_OPERATOR (NULL, CONST_CAST_EXPR, "cc", OVL_OP_FLAG_UNARY)
 DEF_OPERATOR (NULL, STATIC_CAST_EXPR, "sc", OVL_OP_FLAG_UNARY)
+DEF_OPERATOR (NULL, OMP_ARRAYSHAPE_CAST_EXPR, "oc", OVL_OP_FLAG_UNARY)
 DEF_OPERATOR (NULL, SCOPE_REF, "sr", OVL_OP_FLAG_NONE)
 DEF_OPERATOR (NULL, EXPR_PACK_EXPANSION, "sp", OVL_OP_FLAG_NONE)
 DEF_OPERATOR (NULL, UNARY_LEFT_FOLD_EXPR, "fl", OVL_OP_FLAG_NONE)
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index 9a6918af8bb7..994b6d41753f 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -4347,6 +4347,12 @@ cp_parser_new (cp_lexer *lexer)
   /* Disallow OpenMP array sections in expressions.  */
   parser->omp_array_section_p = false;
 
+  /* Disallow OpenMP array-shaping operator in expressions.  */
+  parser->omp_array_shaping_op_p = false;
+
+  /* We don't have an OpenMP array shape here.  */
+  parser->omp_has_array_shape_p = false;
+
   /* Not declaring an implicit function template.  */
   parser->auto_is_implicit_function_template_parm_p = false;
   parser->fully_implicit_function_template_p = false;
@@ -5283,6 +5289,7 @@ cp_parser_statement_expr (cp_parser *parser)
 {
   cp_token_position start = cp_parser_start_tentative_firewall (parser);
   auto oas = make_temp_override (parser->omp_array_section_p, false);
+  auto aso = make_temp_override (parser->omp_array_shaping_op_p, false);
 
   /* Consume the '('.  */
   location_t start_loc = cp_lexer_peek_token (parser->lexer)->location;
@@ -8193,7 +8200,7 @@ cp_parser_postfix_open_square_expression (cp_parser *parser,
       && cp_lexer_next_token_is (parser->lexer, CPP_COLON))
     {
       cp_lexer_consume_token (parser->lexer);
-      tree length = NULL_TREE;
+      tree length = NULL_TREE, stride = NULL_TREE;
       if (cp_lexer_next_token_is_not (parser->lexer, CPP_CLOSE_SQUARE))
 	{
 	  if (cxx_dialect >= cxx23)
@@ -8226,9 +8233,23 @@ cp_parser_postfix_open_square_expression (cp_parser *parser,
 				      /*warn_comma_p=*/warn_comma_subscript);
 	}
 
+      if (cp_lexer_next_token_is (parser->lexer, CPP_COLON))
+	{
+	  cp_lexer_consume_token (parser->lexer);
+	  /* We could check for C++-23 multidimensional/comma-separated
+	     subscripts here, or not bother.  */
+	  if (cp_lexer_next_token_is_not (parser->lexer, CPP_CLOSE_SQUARE))
+	    stride
+	      = cp_parser_expression (parser, NULL, /*cast_p=*/false,
+				      /*decltype_p=*/false,
+				      /*warn_comma_p=*/warn_comma_subscript);
+	}
+
       parser->colon_corrects_to_scope_p = saved_colon_corrects_to_scope_p;
 
-      if (index == error_mark_node || length == error_mark_node)
+      if (index == error_mark_node
+	  || length == error_mark_node
+	  || stride == error_mark_node)
 	{
 	  cp_parser_skip_to_closing_square_bracket (parser);
 	  return error_mark_node;
@@ -8237,7 +8258,7 @@ cp_parser_postfix_open_square_expression (cp_parser *parser,
 	cp_parser_require (parser, CPP_CLOSE_SQUARE, RT_CLOSE_SQUARE);
 
       return grok_omp_array_section (input_location, postfix_expression, index,
-				     length);
+				     length, stride);
     }
 
   parser->colon_corrects_to_scope_p = saved_colon_corrects_to_scope_p;
@@ -8245,11 +8266,23 @@ cp_parser_postfix_open_square_expression (cp_parser *parser,
   /* Look for the closing `]'.  */
   cp_parser_require (parser, CPP_CLOSE_SQUARE, RT_CLOSE_SQUARE);
 
-  /* Build the ARRAY_REF.  */
-  postfix_expression = grok_array_decl (loc, postfix_expression,
-					index, &expression_list,
-					tf_warning_or_error
-					| (decltype_p ? tf_decltype : 0));
+  if (parser->omp_has_array_shape_p
+      && (expression_list.get () == NULL
+	  || vec_safe_length (expression_list) == 1))
+    /* If we have an array-shaping operator, we may not be able to represent
+       a well-formed ARRAY_REF here, because we are coercing the type of the
+       innermost array base and the original type may not be compatible.  Use
+       the OMP_ARRAY_SECTION code instead.  We also want to explicitly avoid
+       creating INDIRECT_REFs for pointer bases, because that can lead to
+       parsing ambiguities (see cp_parser_omp_var_list_no_open).  */
+    return grok_omp_array_section (loc, postfix_expression, index,
+				   size_one_node, NULL_TREE);
+  else
+    /* Build the ARRAY_REF.  */
+    postfix_expression = grok_array_decl (loc, postfix_expression,
+					  index, &expression_list,
+					  tf_warning_or_error
+					  | (decltype_p ? tf_decltype : 0));
 
   /* When not doing offsetof, array references are not permitted in
      constant-expressions.  */
@@ -8571,6 +8604,7 @@ cp_parser_parenthesized_expression_list (cp_parser* parser,
   vec<tree, va_gc> *expression_list;
   bool saved_greater_than_is_operator_p;
   bool saved_omp_array_section_p;
+  bool saved_omp_array_shaping_op_p;
 
   /* Assume all the expressions will be constant.  */
   if (non_constant_p)
@@ -8589,7 +8623,9 @@ cp_parser_parenthesized_expression_list (cp_parser* parser,
   parser->greater_than_is_operator_p = true;
 
   saved_omp_array_section_p = parser->omp_array_section_p;
+  saved_omp_array_shaping_op_p = parser->omp_array_shaping_op_p;
   parser->omp_array_section_p = false;
+  parser->omp_array_shaping_op_p = false;
 
   cp_expr expr (NULL_TREE);
 
@@ -8656,6 +8692,7 @@ cp_parser_parenthesized_expression_list (cp_parser* parser,
 	  parser->greater_than_is_operator_p
 	    = saved_greater_than_is_operator_p;
 	  parser->omp_array_section_p = saved_omp_array_section_p;
+	  parser->omp_array_shaping_op_p = saved_omp_array_shaping_op_p;
 	  return NULL;
 	}
     }
@@ -8663,6 +8700,7 @@ cp_parser_parenthesized_expression_list (cp_parser* parser,
   parser->greater_than_is_operator_p
     = saved_greater_than_is_operator_p;
   parser->omp_array_section_p = saved_omp_array_section_p;
+  parser->omp_array_shaping_op_p = saved_omp_array_shaping_op_p;
 
   return expression_list;
 }
@@ -9931,6 +9969,8 @@ cp_parser_cast_expression (cp_parser *parser, bool address_p, bool cast_p,
       cp_expr expr (NULL_TREE);
       int cast_expression = 0;
       const char *saved_message;
+      auto_vec<cp_expr, 4> omp_shape_dims;
+      bool omp_array_shape_p = false;
 
       /* There's no way to know yet whether or not this is a cast.
 	 For example, `(int (3))' is a unary-expression, while `(int)
@@ -10000,6 +10040,28 @@ cp_parser_cast_expression (cp_parser *parser, bool address_p, bool cast_p,
 	 that the call to cp_parser_error_occurred below returns true.  */
       if (!cast_expression)
 	cp_parser_simulate_error (parser);
+      else if (parser->omp_array_shaping_op_p
+	       && cp_lexer_next_token_is (parser->lexer, CPP_OPEN_SQUARE))
+	{
+	  auto oas = make_temp_override (parser->omp_array_section_p, false);
+	  auto aso = make_temp_override (parser->omp_array_shaping_op_p, false);
+
+	  while (cp_lexer_next_token_is (parser->lexer, CPP_OPEN_SQUARE))
+	    {
+	      cp_lexer_consume_token (parser->lexer);
+	      cp_expr e = cp_parser_expression (parser);
+	      if (e.get_value () == error_mark_node)
+		break;
+	      omp_shape_dims.safe_push (e);
+	      if (!cp_parser_require (parser, CPP_CLOSE_SQUARE,
+				      RT_CLOSE_SQUARE))
+		break;
+	    }
+	  cp_token *close_paren = parens.require_close (parser);
+	  if (close_paren)
+	    close_paren_loc = close_paren->location;
+	  omp_array_shape_p = true;
+	}
       else
 	{
 	  bool saved_in_type_id_in_expr_p = parser->in_type_id_in_expr_p;
@@ -10021,6 +10083,10 @@ cp_parser_cast_expression (cp_parser *parser, bool address_p, bool cast_p,
 	 function returning T.  */
       if (!cp_parser_error_occurred (parser))
 	{
+	  auto aso = make_temp_override (parser->omp_array_shaping_op_p, false);
+	  auto as = make_temp_override (parser->omp_has_array_shape_p,
+					omp_array_shape_p);
+
 	  /* Only commit if the cast-expression doesn't start with
 	     '++', '--', or '[' in C++11.  */
 	  if (cast_expression > 0)
@@ -10034,6 +10100,24 @@ cp_parser_cast_expression (cp_parser *parser, bool address_p, bool cast_p,
 
 	  if (cp_parser_parse_definitely (parser))
 	    {
+	      if (omp_array_shape_p)
+		{
+		  location_t cast_loc = make_location (open_paren_loc,
+						       open_paren_loc,
+						       expr.get_finish ());
+
+		  type = cp_omp_create_arrayshape_type (cast_loc, expr,
+							&omp_shape_dims);
+
+		  /* Things rapidly get worse below if we carry on from here
+		     with an erroneous type...  */
+		  if (error_operand_p (type))
+		    return error_mark_node;
+
+		  return cp_build_omp_arrayshape_cast (cast_loc, type, expr,
+						       tf_warning_or_error);
+		}
+
 	      /* Warn about old-style casts, if so requested.  */
 	      if (warn_old_style_cast
 		  && !in_system_header_at (input_location)
@@ -11148,6 +11232,7 @@ cp_parser_lambda_expression (cp_parser* parser)
     bool auto_is_implicit_function_template_parm_p
         = parser->auto_is_implicit_function_template_parm_p;
     bool saved_omp_array_section_p = parser->omp_array_section_p;
+    bool saved_omp_array_shaping_op_p = parser->omp_array_shaping_op_p;
 
     parser->num_template_parameter_lists = 0;
     parser->in_statement = 0;
@@ -11157,6 +11242,7 @@ cp_parser_lambda_expression (cp_parser* parser)
     parser->implicit_template_scope = 0;
     parser->auto_is_implicit_function_template_parm_p = false;
     parser->omp_array_section_p = false;
+    parser->omp_array_shaping_op_p = false;
 
     /* The body of a lambda in a discarded statement is not discarded.  */
     bool discarded = in_discarded_stmt;
@@ -11208,6 +11294,7 @@ cp_parser_lambda_expression (cp_parser* parser)
     parser->auto_is_implicit_function_template_parm_p
 	= auto_is_implicit_function_template_parm_p;
     parser->omp_array_section_p = saved_omp_array_section_p;
+    parser->omp_array_shaping_op_p = saved_omp_array_shaping_op_p;
   }
 
   /* This field is only used during parsing of the lambda.  */
@@ -25542,6 +25629,7 @@ cp_parser_braced_list (cp_parser* parser, bool* non_constant_p)
   tree initializer;
   location_t start_loc = cp_lexer_peek_token (parser->lexer)->location;
   auto oas = make_temp_override (parser->omp_array_section_p, false);
+  auto aso = make_temp_override (parser->omp_array_shaping_op_p, false);
 
   /* Consume the `{' token.  */
   matching_braces braces;
@@ -37407,11 +37495,11 @@ check_no_duplicate_clause (tree clauses, enum omp_clause_code code,
 
 struct omp_dim
 {
-  tree low_bound, length;
+  tree low_bound, length, stride;
   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) {}
+  omp_dim (tree lb, tree len, tree str, location_t lo, bool nc)
+    : low_bound (lb), length (len), stride (str), loc (lo), no_colon (nc) {}
 };
 
 static tree
@@ -37443,10 +37531,22 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
 		   || kind == OMP_CLAUSE_FROM))
 	{
 	  auto s = make_temp_override (parser->omp_array_section_p, true);
+	  auto o = make_temp_override (parser->omp_array_shaping_op_p,
+				       (kind == OMP_CLAUSE_TO
+					|| kind == OMP_CLAUSE_FROM));
+	  tree reshaped_to = NULL_TREE;
 	  token = cp_lexer_peek_token (parser->lexer);
 	  location_t loc = token->location;
 	  decl = cp_parser_assignment_expression (parser);
 
+	  if ((TREE_CODE (decl) == VIEW_CONVERT_EXPR
+	       && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
+	      || TREE_CODE (decl) == OMP_ARRAYSHAPE_CAST_EXPR)
+	    {
+	      reshaped_to = TREE_TYPE (decl);
+	      decl = TREE_OPERAND (decl, 0);
+	    }
+
 	  /* This code rewrites a parsed expression containing various tree
 	     codes used to represent array accesses into a more uniform nest of
 	     OMP_ARRAY_SECTION nodes before it is processed by
@@ -37457,49 +37557,159 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
 	  dims.truncate (0);
 	  if (TREE_CODE (decl) == OMP_ARRAY_SECTION)
 	    {
+	      size_t sections = 0;
+	      tree orig_decl = decl;
+	      bool update_p = (kind == OMP_CLAUSE_TO
+			       || kind == OMP_CLAUSE_FROM);
+	      bool maybe_ptr_based_noncontig_update = false;
+
+	      while (update_p
+		     && !reshaped_to
+		     && (TREE_CODE (decl) == OMP_ARRAY_SECTION
+			 || TREE_CODE (decl) == ARRAY_REF
+			 || TREE_CODE (decl) == COMPOUND_EXPR))
+		{
+		  if (TREE_CODE (decl) == COMPOUND_EXPR)
+		    decl = TREE_OPERAND (decl, 1);
+		  else
+		    {
+		      if (TREE_CODE (decl) == OMP_ARRAY_SECTION)
+			maybe_ptr_based_noncontig_update = true;
+		      decl = TREE_OPERAND (decl, 0);
+		      sections++;
+		    }
+		}
+
+	      decl = orig_decl;
+
 	      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));
+		  tree stride = TREE_OPERAND (decl, 3);
+		  dims.safe_push (omp_dim (low_bound, length, stride, loc,
+					   false));
 		  decl = TREE_OPERAND (decl, 0);
+		  if (sections > 0)
+		    sections--;
 		}
 
+	      /* The handling of INDIRECT_REF here in the presence of
+		 array-shaping operations is a little tricky.  We need to
+		 avoid treating a pointer dereference as a unit-sized array
+		 section when we have an array shaping operation, because we
+		 don't want an indirection to consume one of the user's
+		 requested array dimensions.  E.g. if we have a
+		 double-indirect pointer like:
+
+		   int **foopp;
+		   #pragma omp target update from(([N][N]) (*foopp)[0:X][0:Y])
+
+		 We don't want to interpret this as:
+
+		   foopp[0:1][0:X][0:Y]
+
+		 else the array shape [N][N] won't match.  Also we can't match
+		 the array sections right-to-left instead, else this:
+
+		   #pragma omp target update from(([N][N]) (*foopp)[0:X])
+
+		 would not copy the dimensions:
+
+		   (*foopp)[0:X][0:N]
+
+		 as required.  So, avoid descending through INDIRECT_REFs if
+		 we have an array-shaping op.
+
+		 If we *don't* have an array-shaping op, but we have a
+		 multiply-indirected pointer and an array section like this:
+
+		   int ***fooppp;
+		   #pragma omp target update from((**fooppp)[0:X:S]
+
+		 also avoid descending through more indirections than we have
+		 array sections, since the noncontiguous update processing code
+		 won't understand them (and doesn't need to traverse them
+		 anyway).  */
+
 	      while (TREE_CODE (decl) == ARRAY_REF
-		     || TREE_CODE (decl) == INDIRECT_REF
+		     || (TREE_CODE (decl) == INDIRECT_REF
+			 && !reshaped_to)
 		     || TREE_CODE (decl) == COMPOUND_EXPR)
 		{
 		  if (REFERENCE_REF_P (decl))
 		    break;
 
+		  if (maybe_ptr_based_noncontig_update && sections == 0)
+		    break;
+
 		  if (TREE_CODE (decl) == COMPOUND_EXPR)
 		    {
 		      decl = TREE_OPERAND (decl, 1);
 		      STRIP_NOPS (decl);
+		      continue;
 		    }
-		  else if (TREE_CODE (decl) == INDIRECT_REF)
+		  else if (TREE_CODE (decl) == INDIRECT_REF
+			   && !reshaped_to)
 		    {
 		      dims.safe_push (omp_dim (integer_zero_node,
-					       integer_one_node, loc, true));
+					       integer_one_node, NULL_TREE, 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));
+		      dims.safe_push (omp_dim (index, integer_one_node,
+					       NULL_TREE, loc, true));
 		      decl = TREE_OPERAND (decl, 0);
+		      if (sections > 0)
+			sections--;
 		    }
 		}
 
+	      if (reshaped_to)
+		{
+		  unsigned reshaped_dims = 0;
+
+		  for (tree t = reshaped_to;
+		       TREE_CODE (t) == ARRAY_TYPE;
+		       t = TREE_TYPE (t))
+		    reshaped_dims++;
+
+		  if (dims.length () > reshaped_dims)
+		    {
+		      error_at (loc, "too many array section specifiers "
+				"for %qT", reshaped_to);
+		      decl = error_mark_node;
+		    }
+		  else
+		    {
+		      /* We have a pointer DECL whose target should be
+			 interpreted as an array with particular dimensions,
+			 not "the pointer itself".  So, add an indirection
+			 here.  */
+		      if (type_dependent_expression_p (decl))
+			decl = build_min_nt_loc (loc, INDIRECT_REF, decl);
+		      else
+			{
+			  /* We're interested in the reference target.  */
+			  decl = convert_from_reference (decl);
+			  decl = cp_build_fold_indirect_ref (decl);
+			}
+		      decl
+			= cp_build_omp_arrayshape_cast (loc, reshaped_to, decl,
+							tf_warning_or_error);
+		    }
+		}
 	      /* Bare references have their own special handling, so remove
 		 the explicit dereference added by convert_from_reference.  */
-	      if (REFERENCE_REF_P (decl))
+	      else if (REFERENCE_REF_P (decl))
 		decl = TREE_OPERAND (decl, 0);
 
 	      for (int i = dims.length () - 1; i >= 0; i--)
 		decl = grok_omp_array_section (loc, decl, dims[i].low_bound,
-					       dims[i].length);
+					       dims[i].length, dims[i].stride);
 	    }
 	  else if (TREE_CODE (decl) == INDIRECT_REF)
 	    {
@@ -37516,7 +37726,7 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
 		   "foo[0:1]".  */
 	      if (!ref_p)
 		decl = grok_omp_array_section (loc, decl, integer_zero_node,
-					       integer_one_node);
+					       integer_one_node, NULL_TREE);
 	    }
 	  else if (TREE_CODE (decl) == ARRAY_REF)
 	    {
@@ -37525,7 +37735,16 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
 	      decl = TREE_OPERAND (decl, 0);
 	      STRIP_NOPS (decl);
 
-	      decl = grok_omp_array_section (loc, decl, idx, integer_one_node);
+	      decl = grok_omp_array_section (loc, decl, idx, integer_one_node,
+					     NULL_TREE);
+	    }
+	  else if (reshaped_to)
+	    {
+	      /* We're copying the whole of a reshaped array, originally a
+		 base pointer.  Rewrite as an array section.  */
+	      tree elems = array_type_nelts_total (reshaped_to);
+	      decl = grok_omp_array_section (loc, decl, size_zero_node, elems,
+					     NULL_TREE);
 	    }
 	  else if (TREE_CODE (decl) == NON_LVALUE_EXPR
 		   || CONVERT_EXPR_P (decl))
@@ -37689,7 +37908,8 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
 		      goto skip_comma;
 		    }
 
-		  dims.safe_push (omp_dim (low_bound, length, loc, no_colon));
+		  dims.safe_push (omp_dim (low_bound, length, NULL_TREE, loc,
+					   no_colon));
 		}
 
 	      if ((kind == OMP_CLAUSE_MAP
@@ -37711,7 +37931,8 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
 		for (unsigned i = 0; i < dims.length (); i++)
 		  decl = build_omp_array_section (input_location, decl,
 						  dims[i].low_bound,
-						  dims[i].length);
+						  dims[i].length,
+						  dims[i].stride);
 	      break;
 	    default:
 	      break;
@@ -37724,6 +37945,8 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
 		  && cp_parser_simulate_error (parser))
 		{
 		depend_lvalue:
+		  auto o = make_temp_override (parser->omp_array_shaping_op_p,
+					       true);
 		  cp_parser_abort_tentative_parse (parser);
 		  decl = cp_parser_assignment_expression (parser, NULL,
 							  false, false);
@@ -45557,8 +45780,38 @@ cp_parser_omp_target_update (cp_parser *parser, cp_token *pragma_tok,
   tree clauses
     = cp_parser_omp_all_clauses (parser, OMP_TARGET_UPDATE_CLAUSE_MASK,
 				 "#pragma omp target update", pragma_tok);
-  if (omp_find_clause (clauses, OMP_CLAUSE_TO) == NULL_TREE
-      && omp_find_clause (clauses, OMP_CLAUSE_FROM) == NULL_TREE)
+  bool to_clause = false, from_clause = false;
+  for (tree c = clauses;
+       c && !to_clause && !from_clause;
+       c = OMP_CLAUSE_CHAIN (c))
+    {
+      switch (OMP_CLAUSE_CODE (c))
+	{
+	case OMP_CLAUSE_TO:
+	  to_clause = true;
+	  break;
+	case OMP_CLAUSE_FROM:
+	  from_clause = true;
+	  break;
+	case OMP_CLAUSE_MAP:
+	  switch (OMP_CLAUSE_MAP_KIND (c))
+	    {
+	    case GOMP_MAP_TO_GRID:
+	      to_clause = true;
+	      break;
+	    case GOMP_MAP_FROM_GRID:
+	      from_clause = true;
+	      break;
+	    default:
+	      ;
+	    }
+	  break;
+	default:
+	  ;
+	}
+    }
+
+  if (!to_clause && !from_clause)
     {
       error_at (pragma_tok->location,
 		"%<#pragma omp target update%> must contain at least one "
diff --git a/gcc/cp/parser.h b/gcc/cp/parser.h
index bb8d7b31db82..0a52f65b9e3d 100644
--- a/gcc/cp/parser.h
+++ b/gcc/cp/parser.h
@@ -410,6 +410,13 @@ struct GTY(()) cp_parser {
   /* TRUE if an OpenMP array section is allowed.  */
   bool omp_array_section_p;
 
+  /* TRUE if an OpenMP array-shaping operator is allowed.  */
+  bool omp_array_shaping_op_p;
+
+  /* TRUE if we are parsing an expression with an OpenMP array-shaping
+     operator.  */
+  bool omp_has_array_shape_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/pt.cc b/gcc/cp/pt.cc
index 80add9d32aa0..bd854bfea7d6 100644
--- a/gcc/cp/pt.cc
+++ b/gcc/cp/pt.cc
@@ -16574,6 +16574,10 @@ tsubst (tree t, tree args, tsubst_flags_t complain, tree in_decl)
 		 member access.  */
 	      id = false;
 	    type = finish_decltype_type (type, id, complain);
+
+	    if (DECLTYPE_FOR_OMP_ARRAYSHAPE_CAST (t)
+		&& TYPE_REF_P (type))
+	      type = TREE_TYPE (type);
 	  }
 	return cp_build_qualified_type (type,
 					cp_type_quals (t)
@@ -17392,6 +17396,7 @@ tsubst_copy (tree t, tree args, tsubst_flags_t complain, tree in_decl)
     case STATIC_CAST_EXPR:
     case DYNAMIC_CAST_EXPR:
     case IMPLICIT_CONV_EXPR:
+    case OMP_ARRAYSHAPE_CAST_EXPR:
     CASE_CONVERT:
       {
 	tsubst_flags_t tcomplain = complain;
@@ -17617,12 +17622,14 @@ tsubst_copy (tree t, tree args, tsubst_flags_t complain, tree in_decl)
     case OMP_ARRAY_SECTION:
       {
 	tree op0 = tsubst_copy (TREE_OPERAND (t, 0), args, complain, in_decl);
-	tree op1 = NULL_TREE, op2 = NULL_TREE;
+	tree op1 = NULL_TREE, op2 = NULL_TREE, op3 = NULL_TREE;
 	if (TREE_OPERAND (t, 1))
 	  op1 = tsubst_copy (TREE_OPERAND (t, 1), args, complain, in_decl);
 	if (TREE_OPERAND (t, 2))
 	  op2 = tsubst_copy (TREE_OPERAND (t, 2), args, complain, in_decl);
-	return build_nt (OMP_ARRAY_SECTION, op0, op1, op2);
+	if (TREE_OPERAND (t, 3))
+	  op3 = tsubst_copy (TREE_OPERAND (t, 3), args, complain, in_decl);
+	return build_nt (OMP_ARRAY_SECTION, op0, op1, op2, op3);
       }
 
     case CALL_EXPR:
@@ -17916,14 +17923,17 @@ tsubst_omp_clause_decl (tree decl, tree args, tsubst_flags_t complain,
 	= tsubst_expr (TREE_OPERAND (decl, 1), args, complain, in_decl);
       tree length = tsubst_expr (TREE_OPERAND (decl, 2), args, complain,
 				 in_decl);
+      tree stride = tsubst_expr (TREE_OPERAND (decl, 3), args, complain,
+				 in_decl);
       tree base = tsubst_omp_clause_decl (TREE_OPERAND (decl, 0), args,
 					   complain, in_decl, NULL);
       if (TREE_OPERAND (decl, 0) == base
 	  && TREE_OPERAND (decl, 1) == low_bound
-	  && TREE_OPERAND (decl, 2) == length)
+	  && TREE_OPERAND (decl, 2) == length
+	  && TREE_OPERAND (decl, 3) == stride)
 	return decl;
-      tree ret = build3 (OMP_ARRAY_SECTION, TREE_TYPE (base), base, low_bound,
-			 length);
+      tree ret = build4 (OMP_ARRAY_SECTION, TREE_TYPE (base), base, low_bound,
+			 length, stride);
       return ret;
     }
   tree ret = tsubst_expr (decl, args, complain, in_decl);
@@ -20524,6 +20534,14 @@ tsubst_copy_and_build (tree t,
 	RETURN (cp_build_bit_cast (EXPR_LOCATION (t), type, op0, complain));
       }
 
+    case OMP_ARRAYSHAPE_CAST_EXPR:
+      {
+	tree type = tsubst (TREE_TYPE (t), args, complain, in_decl);
+	tree op0 = RECUR (TREE_OPERAND (t, 0));
+	RETURN (cp_build_omp_arrayshape_cast (EXPR_LOCATION (t), type, op0,
+					      complain));
+      }
+
     case POSTDECREMENT_EXPR:
     case POSTINCREMENT_EXPR:
       op1 = tsubst_non_call_postfix_expression (TREE_OPERAND (t, 0),
@@ -20690,7 +20708,7 @@ tsubst_copy_and_build (tree t,
     case OMP_ARRAY_SECTION:
       {
 	tree op0 = RECUR (TREE_OPERAND (t, 0));
-	tree op1 = NULL_TREE, op2 = NULL_TREE;
+	tree op1 = NULL_TREE, op2 = NULL_TREE, op3 = NULL_TREE;
 	if (op0 == error_mark_node)
 	  RETURN (error_mark_node);
 	if (TREE_OPERAND (t, 1))
@@ -20705,7 +20723,14 @@ tsubst_copy_and_build (tree t,
 	    if (op2 == error_mark_node)
 	      RETURN (error_mark_node);
 	  }
-	RETURN (build_omp_array_section (EXPR_LOCATION (t), op0, op1, op2));
+	if (TREE_OPERAND (t, 3))
+	  {
+	    op3 = RECUR (TREE_OPERAND (t, 3));
+	    if (op3 == error_mark_node)
+	      RETURN (error_mark_node);
+	  }
+	RETURN (build_omp_array_section (EXPR_LOCATION (t), op0, op1, op2,
+					 op3));
       }
 
     case SIZEOF_EXPR:
diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc
index 1ef00cf9ab3c..c130f6cfaaad 100644
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -5165,9 +5165,9 @@ public:
 static tree
 handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 			     bool &maybe_zero_len, unsigned int &first_non_one,
-			     enum c_omp_region_type ort)
+			     enum c_omp_region_type ort, int *discontiguous)
 {
-  tree ret, low_bound, length, type;
+  tree ret, low_bound, length, stride, type;
   bool openacc = (ort & C_ORT_ACC) != 0;
   if (TREE_CODE (t) != OMP_ARRAY_SECTION)
     {
@@ -5229,18 +5229,26 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
       && TREE_CODE (TREE_OPERAND (t, 0)) == FIELD_DECL)
     TREE_OPERAND (t, 0) = omp_privatize_field (TREE_OPERAND (t, 0), false);
   ret = handle_omp_array_sections_1 (c, TREE_OPERAND (t, 0), types,
-				     maybe_zero_len, first_non_one, ort);
+				     maybe_zero_len, first_non_one, ort,
+				     discontiguous);
   if (ret == error_mark_node || ret == NULL_TREE)
     return ret;
 
-  type = TREE_TYPE (ret);
+  if (TREE_CODE (ret) == OMP_ARRAY_SECTION)
+    type = TREE_TYPE (TREE_TYPE (TREE_OPERAND (ret, 0)));
+  else
+    type = TREE_TYPE (ret);
   low_bound = TREE_OPERAND (t, 1);
   length = TREE_OPERAND (t, 2);
+  stride = TREE_OPERAND (t, 3);
   if ((low_bound && type_dependent_expression_p (low_bound))
-      || (length && type_dependent_expression_p (length)))
+      || (length && type_dependent_expression_p (length))
+      || (stride && type_dependent_expression_p (stride)))
     return NULL_TREE;
 
-  if (low_bound == error_mark_node || length == error_mark_node)
+  if (low_bound == error_mark_node
+      || length == error_mark_node
+      || stride == error_mark_node)
     return error_mark_node;
 
   if (low_bound && !INTEGRAL_TYPE_P (TREE_TYPE (low_bound)))
@@ -5257,15 +5265,26 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 		length);
       return error_mark_node;
     }
+  if (stride && !INTEGRAL_TYPE_P (TREE_TYPE (stride)))
+    {
+      error_at (OMP_CLAUSE_LOCATION (c),
+		"stride %qE of array section does not have integral type",
+		stride);
+      return error_mark_node;
+    }
   if (low_bound)
     low_bound = mark_rvalue_use (low_bound);
   if (length)
     length = mark_rvalue_use (length);
+  if (stride)
+    stride = mark_rvalue_use (stride);
   /* We need to reduce to real constant-values for checks below.  */
   if (length)
     length = fold_simple (length);
   if (low_bound)
     low_bound = fold_simple (low_bound);
+  if (stride)
+    stride = fold_simple (stride);
   if (low_bound
       && TREE_CODE (low_bound) == INTEGER_CST
       && TYPE_PRECISION (TREE_TYPE (low_bound))
@@ -5276,9 +5295,15 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
       && TYPE_PRECISION (TREE_TYPE (length))
 	 > TYPE_PRECISION (sizetype))
     length = fold_convert (sizetype, length);
+  if (stride
+      && TREE_CODE (stride) == INTEGER_CST
+      && TYPE_PRECISION (TREE_TYPE (stride))
+	 > TYPE_PRECISION (sizetype))
+    stride = fold_convert (sizetype, stride);
   if (low_bound == NULL_TREE)
     low_bound = integer_zero_node;
-
+  if (stride == NULL_TREE)
+    stride = size_one_node;
   if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
       && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
 	  || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))
@@ -5397,12 +5422,29 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 	    }
 	  if (length && TREE_CODE (length) == INTEGER_CST)
 	    {
-	      if (tree_int_cst_lt (size, length))
+	      tree slength = length;
+	      if (stride && TREE_CODE (stride) == INTEGER_CST)
 		{
-		  error_at (OMP_CLAUSE_LOCATION (c),
-			    "length %qE above array section size "
-			    "in %qs clause", length,
-			    omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+		  slength = size_binop (MULT_EXPR,
+					fold_convert (sizetype, length),
+					fold_convert (sizetype, stride));
+		  slength = size_binop (MINUS_EXPR,
+					  slength,
+					  fold_convert (sizetype, stride));
+		  slength = size_binop (PLUS_EXPR, slength, size_one_node);
+		}
+	      if (tree_int_cst_lt (size, slength))
+		{
+		  if (stride)
+		    error_at (OMP_CLAUSE_LOCATION (c),
+			      "length %qE with stride %qE above array "
+			      "section size in %qs clause", length, stride,
+			      omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+		  else
+		    error_at (OMP_CLAUSE_LOCATION (c),
+			      "length %qE above array section size "
+			      "in %qs clause", length,
+			      omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
 		  return error_mark_node;
 		}
 	      if (TREE_CODE (low_bound) == INTEGER_CST)
@@ -5410,7 +5452,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 		  tree lbpluslen
 		    = size_binop (PLUS_EXPR,
 				  fold_convert (sizetype, low_bound),
-				  fold_convert (sizetype, length));
+				  fold_convert (sizetype, slength));
 		  if (TREE_CODE (lbpluslen) == INTEGER_CST
 		      && tree_int_cst_lt (size, lbpluslen))
 		    {
@@ -5480,12 +5522,20 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 	       d = TREE_OPERAND (d, 0))
 	    {
 	      tree d_length = TREE_OPERAND (d, 2);
-	      if (d_length == NULL_TREE || !integer_onep (d_length))
+	      tree d_stride = TREE_OPERAND (d, 3);
+	      if (d_length == NULL_TREE
+		  || !integer_onep (d_length)
+		  || (d_stride && !integer_onep (d_stride)))
 		{
-		  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 (discontiguous && *discontiguous)
+		    *discontiguous = 2;
+		  else
+		    {
+		      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;
+		    }
 		}
 	    }
 	}
@@ -5497,7 +5547,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
       return error_mark_node;
     }
   if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND)
-    types.safe_push (TREE_TYPE (ret));
+    types.safe_push (type);
   /* We will need to evaluate lb more than once.  */
   tree lb = cp_save_expr (low_bound);
   if (lb != low_bound)
@@ -5516,15 +5566,45 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 		      OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION
 		      || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IN_REDUCTION
 		      || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TASK_REDUCTION);
-  ret = grok_array_decl (OMP_CLAUSE_LOCATION (c), ret, low_bound, NULL,
-			 tf_warning_or_error);
+  /* NOTE: Stride/length are discarded for affinity/depend here.  */
+  if (discontiguous
+      && *discontiguous
+      && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_AFFINITY
+      && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND)
+    ret = grok_omp_array_section (OMP_CLAUSE_LOCATION (c), ret, low_bound,
+				  length, stride);
+  else
+    ret = grok_array_decl (OMP_CLAUSE_LOCATION (c), ret, low_bound, NULL,
+			   tf_warning_or_error);
   return ret;
 }
 
-/* Handle array sections for clause C.  */
+/* We built a reference to an array section, but it turns out we only need a
+   set of ARRAY_REFs to the lower bound.  Rewrite the node.  */
+
+static tree
+omp_array_section_low_bound (location_t loc, tree node)
+{
+  if (TREE_CODE (node) == OMP_ARRAY_SECTION)
+    {
+      tree low_bound = TREE_OPERAND (node, 1);
+      tree ret
+	= omp_array_section_low_bound (loc, TREE_OPERAND (node, 0));
+      return grok_array_decl (loc, ret, low_bound, NULL, tf_warning_or_error);
+    }
+
+  return node;
+}
+
+/* Handle array sections for clause C.  On entry *DISCONTIGUOUS is 0 if array
+   section must be contiguous, 1 if it can be discontiguous, and in the latter
+   case it is set to 2 on exit if it is determined to be discontiguous during
+   the function's execution.  PC points to the clause to be processed, and
+   *PNEXT to the last mapping node created, if passed as non-NULL.  */
 
 static bool
-handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort)
+handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort,
+			   int *discontiguous)
 {
   tree c = *pc;
   bool maybe_zero_len = false;
@@ -5539,7 +5619,7 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort)
     tp = &TREE_VALUE (*tp);
   tree first = handle_omp_array_sections_1 (c, *tp, types,
 					    maybe_zero_len, first_non_one,
-					    ort);
+					    ort, discontiguous);
   if (first == error_mark_node)
     return true;
   if (first == NULL_TREE)
@@ -5580,6 +5660,8 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort)
       if (processing_template_decl && maybe_zero_len)
 	return false;
 
+      bool higher_discontiguous = false;
+
       for (i = num, t = OMP_CLAUSE_DECL (c); i > 0;
 	   t = TREE_OPERAND (t, 0))
 	{
@@ -5587,6 +5669,7 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort)
 
 	  tree low_bound = TREE_OPERAND (t, 1);
 	  tree length = TREE_OPERAND (t, 2);
+	  tree stride = TREE_OPERAND (t, 3);
 
 	  i--;
 	  if (low_bound
@@ -5599,12 +5682,56 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort)
 	      && TYPE_PRECISION (TREE_TYPE (length))
 		 > TYPE_PRECISION (sizetype))
 	    length = fold_convert (sizetype, length);
+	  if (stride
+	      && TREE_CODE (stride) == INTEGER_CST
+	      && TYPE_PRECISION (TREE_TYPE (stride))
+		 > TYPE_PRECISION (sizetype))
+	    stride = fold_convert (sizetype, stride);
 	  if (low_bound == NULL_TREE)
 	    low_bound = integer_zero_node;
+	  if (stride == NULL_TREE)
+	    stride = size_one_node;
+	  if (discontiguous && *discontiguous)
+	    {
+	      /* This condition is similar to the error check below, but
+		 whereas that checks for a definitely-discontiguous array
+		 section in order to report an error (where such a section is
+		 illegal), here we instead need to know if the array section
+		 *may be* discontiguous so we can handle that case
+		 appropriately (i.e. for rectangular "target update"
+		 operations).  */
+	      bool full_span = false;
+	      if (length != NULL_TREE
+		  && TREE_CODE (length) == INTEGER_CST
+		  && TREE_CODE (types[i]) == ARRAY_TYPE
+		  && TYPE_DOMAIN (types[i])
+		  && TYPE_MAX_VALUE (TYPE_DOMAIN (types[i]))
+		  && TREE_CODE (TYPE_MAX_VALUE (TYPE_DOMAIN (types[i])))
+		     == INTEGER_CST)
+		{
+		  tree size;
+		  size = size_binop (PLUS_EXPR,
+				     TYPE_MAX_VALUE (TYPE_DOMAIN (types[i])),
+				     size_one_node);
+		  if (tree_int_cst_equal (length, size))
+		    full_span = true;
+		}
+
+	      if (!integer_onep (stride)
+		  || (higher_discontiguous
+		      && (!integer_zerop (low_bound)
+			  || !full_span)))
+		*discontiguous = 2;
+
+	      if (!integer_onep (stride)
+		  || !integer_zerop (low_bound)
+		  || !full_span)
+		higher_discontiguous = true;
+	    }
 	  if (!maybe_zero_len && i > first_non_one)
 	    {
 	      if (integer_nonzerop (low_bound))
-		goto do_warn_noncontiguous;
+		goto is_noncontiguous;
 	      if (length != NULL_TREE
 		  && TREE_CODE (length) == INTEGER_CST
 		  && TYPE_DOMAIN (types[i])
@@ -5618,12 +5745,17 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort)
 				     size_one_node);
 		  if (!tree_int_cst_equal (length, size))
 		    {
-		     do_warn_noncontiguous:
-		      error_at (OMP_CLAUSE_LOCATION (c),
-				"array section is not contiguous in %qs "
-				"clause",
-				omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
-		      return true;
+		     is_noncontiguous:
+		      if (discontiguous && *discontiguous)
+			*discontiguous = 2;
+		      else
+			{
+			  error_at (OMP_CLAUSE_LOCATION (c),
+				    "array section is not contiguous in %qs "
+				    "clause",
+				    omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+			  return true;
+			}
 		    }
 		}
 	      if (!processing_template_decl
@@ -5732,6 +5864,9 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort)
 	      OMP_CLAUSE_DECL (c) = t;
 	      return false;
 	    }
+	  if (discontiguous && *discontiguous != 2)
+	    first = omp_array_section_low_bound (OMP_CLAUSE_LOCATION (c),
+						 first);
 	  OMP_CLAUSE_DECL (c) = first;
 	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
 	    return false;
@@ -5739,9 +5874,6 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort)
 	  if (TREE_CODE (t) == FIELD_DECL)
 	    t = finish_non_static_data_member (t, NULL_TREE, NULL_TREE);
 
-	  if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
-	    return false;
-
 	  if (TREE_CODE (first) == INDIRECT_REF)
 	    {
 	      /* Detect and skip adding extra nodes for pointer-to-member
@@ -5768,6 +5900,10 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort)
 		}
 	    }
 
+	  if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+	      && !(discontiguous && *discontiguous == 2))
+	    return false;
+
 	  /* FIRST represents the first item of data that we are mapping.
 	     E.g. if we're mapping an array, FIRST might resemble
 	     "foo.bar.myarray[0]".  */
@@ -5786,7 +5922,8 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort)
 
 	      c = *pc;
 
-	      if (ai.maybe_zero_length_array_section (c))
+	      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+		  && ai.maybe_zero_length_array_section (c))
 		OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1;
 
 	      /* !!! If we're accessing a base decl via chained access
@@ -6925,7 +7062,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	  t = OMP_CLAUSE_DECL (c);
 	  if (TREE_CODE (t) == OMP_ARRAY_SECTION)
 	    {
-	      if (handle_omp_array_sections (pc, NULL, ort))
+	      if (handle_omp_array_sections (pc, NULL, ort, NULL))
 		{
 		  remove = true;
 		  break;
@@ -7965,7 +8102,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 
 	  if (TREE_CODE (t) == OMP_ARRAY_SECTION)
 	    {
-	      if (handle_omp_array_sections (pc, NULL, ort))
+	      int discontiguous = 1;
+	      if (handle_omp_array_sections (pc, NULL, ort, &discontiguous))
 		remove = true;
 	      else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND
 		       && (OMP_CLAUSE_DEPEND_KIND (c)
@@ -8120,6 +8258,9 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	      remove = true;
 	      break;
 	    }
+	  if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_DIM
+	      || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_STRIDE)
+	    break;
 	  /* FALLTHRU */
 	case OMP_CLAUSE_TO:
 	case OMP_CLAUSE_FROM:
@@ -8134,8 +8275,11 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		grp_start_p = pc;
 		grp_sentinel = OMP_CLAUSE_CHAIN (c);
 
+		int discontiguous
+		  = (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO
+		     || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM);
 		tree *pnext = NULL;
-		if (handle_omp_array_sections (pc, &pnext, ort))
+		if (handle_omp_array_sections (pc, &pnext, ort, &discontiguous))
 		  remove = true;
 		else
 		  {
@@ -8723,7 +8867,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	  t = OMP_CLAUSE_DECL (c);
 	  if (TREE_CODE (t) == OMP_ARRAY_SECTION)
 	    {
-	      if (handle_omp_array_sections (pc, NULL, ort))
+	      if (handle_omp_array_sections (pc, NULL, ort, NULL))
 		remove = true;
 	      else
 		{
@@ -12719,4 +12863,43 @@ cp_build_bit_cast (location_t loc, tree type, tree arg,
   return ret;
 }
 
+/* Build an OpenMP array-shape cast of ARG to TYPE.  */
+
+tree
+cp_build_omp_arrayshape_cast (location_t loc, tree type, tree arg,
+			      tsubst_flags_t complain)
+{
+  if (error_operand_p (type))
+    return error_mark_node;
+
+  if (!dependent_type_p (type)
+      && !complete_type_or_maybe_complain (type, NULL_TREE, complain))
+    return error_mark_node;
+
+  if (error_operand_p (arg))
+    return error_mark_node;
+
+  if (!type_dependent_expression_p (arg) && !dependent_type_p (type))
+    {
+      if (!trivially_copyable_p (TREE_TYPE (arg)))
+	{
+	  error_at (cp_expr_loc_or_loc (arg, loc),
+		    "OpenMP array shape source type %qT "
+		    "is not trivially copyable", TREE_TYPE (arg));
+	  return error_mark_node;
+	}
+
+      /* A pointer to multi-dimensional array conversion isn't normally
+	 allowed, but we force it here for array shape operators by creating
+	 the node directly.  We also want to avoid any overloaded conversions
+	 the user might have defined, not that there are likely to be any.  */
+      return build1_loc (loc, VIEW_CONVERT_EXPR, type, arg);
+    }
+
+  tree ret = build_min (OMP_ARRAYSHAPE_CAST_EXPR, type, arg);
+  SET_EXPR_LOCATION (ret, loc);
+
+  return ret;
+}
+
 #include "gt-cp-semantics.h"
diff --git a/gcc/cp/typeck.cc b/gcc/cp/typeck.cc
index 73890c0c27bc..195c3dec1154 100644
--- a/gcc/cp/typeck.cc
+++ b/gcc/cp/typeck.cc
@@ -1625,6 +1625,9 @@ structural_comptypes (tree t1, tree t2, int strict)
 	return false;
       if (DECLTYPE_FOR_LAMBDA_PROXY (t1) != DECLTYPE_FOR_LAMBDA_PROXY (t2))
 	return false;
+      if (DECLTYPE_FOR_OMP_ARRAYSHAPE_CAST (t1)
+	  != DECLTYPE_FOR_OMP_ARRAYSHAPE_CAST (t2))
+	return false;
       if (!cp_tree_equal (DECLTYPE_TYPE_EXPR (t1), DECLTYPE_TYPE_EXPR (t2)))
         return false;
       break;
@@ -4793,7 +4796,7 @@ build_x_array_ref (location_t loc, tree arg1, tree arg2,
 
 tree
 build_omp_array_section (location_t loc, tree array_expr, tree index,
-			 tree length)
+			 tree length, tree stride)
 {
   tree idxtype;
 
@@ -4832,8 +4835,8 @@ build_omp_array_section (location_t loc, tree array_expr, tree index,
   else
     sectype = build_array_type (eltype, idxtype);
 
-  return build3_loc (loc, OMP_ARRAY_SECTION, sectype, array_expr, index,
-		     length);
+  return build4_loc (loc, OMP_ARRAY_SECTION, sectype, array_expr, index,
+		     length, stride);
 }
 
 /* Return whether OP is an expression of enum type cast to integer
@@ -8150,6 +8153,9 @@ check_for_casting_away_constness (location_t loc, tree src_type,
 		  src_type, dest_type);
       return true;
 
+    case OMP_ARRAYSHAPE_CAST_EXPR:
+      return true;
+
     default:
       gcc_unreachable();
     }
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index ae2fbc65c690..fd580ccb8a4c 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -9203,6 +9203,19 @@ omp_group_last (tree *start_p)
 	grp_last_p = &OMP_CLAUSE_CHAIN (c);
       break;
 
+    case GOMP_MAP_TO_GRID:
+    case GOMP_MAP_FROM_GRID:
+      while (nc
+	     && OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP
+	     && (OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_DIM
+		 || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_STRIDE))
+	{
+	  grp_last_p = &OMP_CLAUSE_CHAIN (c);
+	  c = nc;
+	    nc = OMP_CLAUSE_CHAIN (c);
+	}
+      break;
+
     case GOMP_MAP_STRUCT:
     case GOMP_MAP_STRUCT_UNORD:
       {
@@ -9346,6 +9359,10 @@ omp_group_base (omp_mapping_group *grp, unsigned int *chained,
 	internal_error ("unexpected mapping node");
       return error_mark_node;
 
+    case GOMP_MAP_TO_GRID:
+    case GOMP_MAP_FROM_GRID:
+      return *grp->grp_start;
+
     case GOMP_MAP_ATTACH:
     case GOMP_MAP_DETACH:
       node = OMP_CLAUSE_CHAIN (node);
@@ -14047,7 +14064,9 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
 	    }
 	  if (remove)
 	    break;
-	  if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
+	  if (OMP_CLAUSE_SIZE (c) == NULL_TREE
+	      && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_GRID_DIM
+	      && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_GRID_STRIDE)
 	    OMP_CLAUSE_SIZE (c) = DECL_P (decl) ? DECL_SIZE_UNIT (decl)
 				  : TYPE_SIZE_UNIT (TREE_TYPE (decl));
 	  gimplify_omp_ctxp = ctx->outer_context;
@@ -14155,6 +14174,20 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
 				 is_gimple_lvalue, fb_lvalue) == GS_ERROR)
 		remove = true;
 	    }
+	  else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_DIM
+		   || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_STRIDE)
+	    {
+	      /* The OMP_CLAUSE_DECL for GRID_DIM/GRID_STRIDE isn't necessarily
+		 an lvalue -- e.g. it might be a constant.  So handle it
+		 specially here.  */
+	      if (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p, NULL,
+				 is_gimple_val, fb_rvalue) == GS_ERROR)
+		{
+		  gimplify_omp_ctxp = ctx;
+		  remove = true;
+		}
+	      break;
+	    }
 	  else if (!DECL_P (decl))
 	    {
 	      if ((ctx->region_type & ORT_TARGET) != 0
@@ -14247,8 +14280,13 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
 
 	      gimplify_omp_ctxp = ctx->outer_context;
 	      if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue,
-				 fb_lvalue) == GS_ERROR)
-		remove = true;
+				 fb_lvalue | fb_mayfail) == GS_ERROR)
+		{
+		  sorry_at (OMP_CLAUSE_LOCATION (c),
+			    "unsupported map expression %qE",
+			    OMP_CLAUSE_DECL (c));
+		  remove = true;
+		}
 	      gimplify_omp_ctxp = ctx;
 
 	      break;
diff --git a/gcc/omp-general.cc b/gcc/omp-general.cc
index caddac939cd9..1f6a4b23baa7 100644
--- a/gcc/omp-general.cc
+++ b/gcc/omp-general.cc
@@ -3174,6 +3174,32 @@ omp_parse_pointer (tree *expr0, bool *has_offset)
   return false;
 }
 
+static bool
+omp_parse_noncontiguous_array (tree *expr0)
+{
+  tree expr = *expr0;
+  bool noncontig = false;
+
+  while (TREE_CODE (expr) == OMP_ARRAY_SECTION
+	 || TREE_CODE (expr) == ARRAY_REF)
+    {
+      /* Contiguous arrays use ARRAY_REF.  By the time we reach here,
+	 OMP_ARRAY_SECTION is only used for noncontiguous arrays.  */
+      if (TREE_CODE (expr) == OMP_ARRAY_SECTION)
+	noncontig = true;
+
+      expr = TREE_OPERAND (expr, 0);
+    }
+
+  if (noncontig)
+    {
+      *expr0 = expr;
+      return true;
+    }
+
+  return false;
+}
+
 static bool
 omp_parse_access_method (tree *expr0, enum access_method_kinds *kind)
 {
@@ -3182,6 +3208,13 @@ omp_parse_access_method (tree *expr0, enum access_method_kinds *kind)
 
   if (omp_parse_ref (&expr))
     *kind = ACCESS_REF;
+  else if (omp_parse_noncontiguous_array (&expr))
+    {
+      if (omp_parse_ref (&expr))
+	*kind = ACCESS_NONCONTIG_REF_TO_ARRAY;
+      else
+	*kind = ACCESS_NONCONTIG_ARRAY;
+    }
   else if (omp_parse_pointer (&expr, &has_offset))
     {
       if (omp_parse_ref (&expr))
@@ -3255,6 +3288,14 @@ omp_parse_structure_base (vec<omp_addr_token *> &addr_tokens,
       return true;
     }
 
+  if (TREE_CODE (expr) == VIEW_CONVERT_EXPR
+      && TREE_CODE (TREE_TYPE (expr)) == ARRAY_TYPE)
+    {
+      *kind = BASE_DECL;
+      *expr0 = TREE_OPERAND (expr, 0);
+      return true;
+    }
+
   *kind = BASE_ARBITRARY_EXPR;
   *expr0 = expr;
   return true;
@@ -3404,6 +3445,12 @@ debug_omp_tokenized_addr (vec<omp_addr_token *> &addr_tokens,
 	    case ACCESS_INDEXED_REF_TO_ARRAY:
 	      fputs ("access_indexed_ref_to_array", stderr);
 	      break;
+	    case ACCESS_NONCONTIG_ARRAY:
+	      fputs ("access_noncontig_array", stderr);
+	      break;
+	    case ACCESS_NONCONTIG_REF_TO_ARRAY:
+	      fputs ("access_noncontig_ref_to_array", stderr);
+	      break;
 	    }
 	  break;
 	case ARRAY_BASE:
diff --git a/gcc/omp-general.h b/gcc/omp-general.h
index 0398a9e95e4a..3fa20d26bbb2 100644
--- a/gcc/omp-general.h
+++ b/gcc/omp-general.h
@@ -254,7 +254,9 @@ enum access_method_kinds
   ACCESS_POINTER_OFFSET,
   ACCESS_REF_TO_POINTER_OFFSET,
   ACCESS_INDEXED_ARRAY,
-  ACCESS_INDEXED_REF_TO_ARRAY
+  ACCESS_INDEXED_REF_TO_ARRAY,
+  ACCESS_NONCONTIG_ARRAY,
+  ACCESS_NONCONTIG_REF_TO_ARRAY
 };
 
 /* These are the kinds that a STRUCTURE_BASE or ARRAY_BASE (except
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index bc3dc5c9fb34..6586a48ca789 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -1141,6 +1141,55 @@ fixup_child_record_type (omp_context *ctx)
     = build_qualified_type (build_reference_type (type), TYPE_QUAL_RESTRICT);
 }
 
+/* Build record type for noncontiguous target update operations.  Must be kept
+   in sync with libgomp/libgomp.h omp_noncontig_array_desc.  */
+
+static tree
+omp_noncontig_descriptor_type (location_t loc)
+{
+  static tree cached = NULL_TREE;
+
+  if (cached)
+    return cached;
+
+  tree t = make_node (RECORD_TYPE);
+
+  tree fields = build_decl (loc, FIELD_DECL, get_identifier ("__ndims"),
+			    size_type_node);
+
+  tree field = build_decl (loc, FIELD_DECL, get_identifier ("__elemsize"),
+			   size_type_node);
+  TREE_CHAIN (field) = fields;
+  fields = field;
+
+  tree ptr_size_type = build_pointer_type (size_type_node);
+
+  field = build_decl (loc, FIELD_DECL, get_identifier ("__dim"), ptr_size_type);
+  TREE_CHAIN (field) = fields;
+  fields = field;
+
+  field = build_decl (loc, FIELD_DECL, get_identifier ("__index"),
+		      ptr_size_type);
+  TREE_CHAIN (field) = fields;
+  fields = field;
+
+  field = build_decl (loc, FIELD_DECL, get_identifier ("__length"),
+		      ptr_size_type);
+  TREE_CHAIN (field) = fields;
+  fields = field;
+
+  field = build_decl (loc, FIELD_DECL, get_identifier ("__stride"),
+		      ptr_size_type);
+  TREE_CHAIN (field) = fields;
+  fields = field;
+
+  finish_builtin_struct (t, "__omp_noncontig_desc_type", fields, ptr_type_node);
+
+  cached = t;
+
+  return t;
+}
+
 /* Instantiate decls as necessary in CTX to satisfy the data sharing
    specified by CLAUSES.  */
 
@@ -1657,7 +1706,74 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	      install_var_local (decl, ctx);
 	      break;
 	    }
-	  if (DECL_P (decl))
+	  else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+		   && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_GRID
+		       || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FROM_GRID))
+	    {
+	      tree desc_type = omp_noncontig_descriptor_type (UNKNOWN_LOCATION);
+
+	      tree bare = decl;
+	      if (TREE_CODE (bare) == VIEW_CONVERT_EXPR)
+		bare = TREE_OPERAND (bare, 0);
+
+	      const char *desc_name = ".omp_noncontig_desc";
+	      /* Try (but not too hard) to make a friendly name for the
+		 descriptor.  */
+	      if (DECL_P (bare))
+		desc_name = ACONCAT ((".omp_nc_desc_",
+				      IDENTIFIER_POINTER (DECL_NAME (bare)),
+				      NULL));
+	      tree desc = create_tmp_var (desc_type, desc_name);
+	      DECL_NAMELESS (desc) = 1;
+	      TREE_ADDRESSABLE (desc) = 1;
+
+	      /* Adjust DECL so it refers to the first element of the array:
+		 either by indirecting a pointer, or by selecting the zero'th
+		 index of each dimension of an array.  (We don't have a "bias"
+		 as such for this type of noncontiguous update operation, just
+		 the volume specified in the descriptor we build in
+		 lower_omp_target.)  */
+
+	      if (TREE_CODE (TREE_TYPE (decl)) == POINTER_TYPE)
+		{
+		  decl = build_fold_indirect_ref (decl);
+		  OMP_CLAUSE_DECL (c) = decl;
+		}
+
+	      tree field
+		= build_decl (OMP_CLAUSE_LOCATION (c), FIELD_DECL, NULL_TREE,
+			      ptr_type_node);
+	      SET_DECL_ALIGN (field, TYPE_ALIGN (ptr_type_node));
+	      insert_field_into_struct (ctx->record_type, field);
+	      splay_tree_insert (ctx->field_map, (splay_tree_key) decl,
+				 (splay_tree_value) field);
+
+	      tree dn = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+					  OMP_CLAUSE_MAP);
+	      OMP_CLAUSE_SET_MAP_KIND (dn, GOMP_MAP_TO_PSET);
+	      OMP_CLAUSE_DECL (dn) = desc;
+	      OMP_CLAUSE_SIZE (dn) = TYPE_SIZE_UNIT (desc_type);
+
+	      OMP_CLAUSE_CHAIN (dn) = OMP_CLAUSE_CHAIN (c);
+	      OMP_CLAUSE_CHAIN (c) = dn;
+
+	      field = build_decl (OMP_CLAUSE_LOCATION (c), FIELD_DECL,
+				  NULL_TREE, ptr_type_node);
+	      SET_DECL_ALIGN (field, TYPE_ALIGN (ptr_type_node));
+	      insert_field_into_struct (ctx->record_type, field);
+	      splay_tree_insert (ctx->field_map, (splay_tree_key) desc,
+				 (splay_tree_value) field);
+
+	      c = dn;
+	      tree nc;
+
+	      while ((nc = OMP_CLAUSE_CHAIN (c))
+		     && OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP
+		     && (OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_DIM
+			 || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_STRIDE))
+		c = nc;
+	    }
+	  else if (DECL_P (decl))
 	    {
 	      if (DECL_SIZE (decl)
 		  && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
@@ -1890,6 +2006,11 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	      && is_omp_target (ctx->stmt)
 	      && !is_gimple_omp_offloaded (ctx->stmt))
 	    break;
+	  if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_GRID
+	      || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FROM_GRID
+	      || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_DIM
+	      || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_STRIDE)
+	    break;
 	  if (DECL_P (decl))
 	    {
 	      if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
@@ -12809,6 +12930,10 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	  case GOMP_MAP_DETACH:
 	  case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
 	  case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION:
+	  case GOMP_MAP_TO_GRID:
+	  case GOMP_MAP_FROM_GRID:
+	  case GOMP_MAP_GRID_DIM:
+	  case GOMP_MAP_GRID_STRIDE:
 	    break;
 	  case GOMP_MAP_IF_PRESENT:
 	  case GOMP_MAP_FORCE_ALLOC:
@@ -12826,6 +12951,20 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	    gcc_unreachable ();
 	  }
 #endif
+	if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_GRID
+	    || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FROM_GRID)
+	  {
+	    tree nc = OMP_CLAUSE_CHAIN (c);
+	    gcc_assert (OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP
+			&& OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_TO_PSET);
+	    c = nc;
+	    while ((nc = OMP_CLAUSE_CHAIN (c))
+		   && (OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_DIM
+		       || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_STRIDE))
+	      c = nc;
+	    map_cnt += 2;
+	    continue;
+	  }
 	  /* FALLTHRU */
       case OMP_CLAUSE_TO:
       case OMP_CLAUSE_FROM:
@@ -13117,7 +13256,267 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 		    || (OMP_CLAUSE_MAP_KIND (c)
 			== GOMP_MAP_FIRSTPRIVATE_REFERENCE)))
 	      break;
-	    if (!DECL_P (ovar))
+	    if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+		&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_GRID
+		    || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FROM_GRID))
+	      {
+		tree decl = OMP_CLAUSE_DECL (c);
+		tree dn = OMP_CLAUSE_CHAIN (c);
+		gcc_assert (OMP_CLAUSE_CODE (dn) == OMP_CLAUSE_MAP
+			    && OMP_CLAUSE_MAP_KIND (dn) == GOMP_MAP_TO_PSET);
+		tree desc = OMP_CLAUSE_DECL (dn);
+
+		tree oc, elsize = OMP_CLAUSE_SIZE (c);
+		tree type = TREE_TYPE (decl);
+		int i, dims = 0;
+		auto_vec<tree> tdims;
+		bool pointer_based = false, handled_pointer_section = false;
+		tree arrsize = fold_convert (sizetype, elsize);
+
+		/* Allow a single (maybe strided) array section if we have a
+		   pointer base.  */
+		if (TREE_CODE (decl) == INDIRECT_REF
+		    && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
+			== POINTER_TYPE))
+		  {
+		    pointer_based = true;
+		    dims = 1;
+		  }
+		else
+		  for (tree itype = type;
+		       TREE_CODE (itype) == ARRAY_TYPE;
+		       itype = TREE_TYPE (itype))
+		    {
+		      tdims.safe_push (itype);
+		      dims++;
+		    }
+
+		int tdim = tdims.length () - 1;
+
+		vec<constructor_elt, va_gc> *vdim;
+		vec<constructor_elt, va_gc> *vindex;
+		vec<constructor_elt, va_gc> *vlen;
+		vec<constructor_elt, va_gc> *vstride;
+		vec_alloc (vdim, dims);
+		vec_alloc (vindex, dims);
+		vec_alloc (vlen, dims);
+		vec_alloc (vstride, dims);
+
+		tree size_arr_type
+		  = build_array_type_nelts (size_type_node, dims);
+
+		tree dim_tmp = create_tmp_var (size_arr_type, ".omp_dim");
+		DECL_NAMELESS (dim_tmp) = 1;
+		TREE_ADDRESSABLE (dim_tmp) = 1;
+		TREE_STATIC (dim_tmp) = 1;
+		tree index_tmp = create_tmp_var (size_arr_type, ".omp_index");
+		DECL_NAMELESS (index_tmp) = 1;
+		TREE_ADDRESSABLE (index_tmp) = 1;
+		TREE_STATIC (index_tmp) = 1;
+		tree len_tmp = create_tmp_var (size_arr_type, ".omp_len");
+		DECL_NAMELESS (len_tmp) = 1;
+		TREE_ADDRESSABLE (len_tmp) = 1;
+		TREE_STATIC (len_tmp) = 1;
+		tree stride_tmp = create_tmp_var (size_arr_type, ".omp_stride");
+		DECL_NAMELESS (stride_tmp) = 1;
+		TREE_ADDRESSABLE (stride_tmp) = 1;
+		TREE_STATIC (stride_tmp) = 1;
+
+		oc = c;
+		c = dn;
+
+		for (i = 0; i < dims; i++)
+		  {
+		    nc = OMP_CLAUSE_CHAIN (c);
+		    tree dim = NULL_TREE, index = NULL_TREE, len = NULL_TREE,
+			 stride = size_one_node;
+
+		    if (OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP
+			&& OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_DIM)
+		      {
+			index = OMP_CLAUSE_DECL (nc);
+			len = OMP_CLAUSE_SIZE (nc);
+
+			index = fold_convert (sizetype, index);
+			len = fold_convert (sizetype, len);
+
+			tree nc2 = OMP_CLAUSE_CHAIN (nc);
+			if (nc2
+			    && OMP_CLAUSE_CODE (nc2) == OMP_CLAUSE_MAP
+			    && (OMP_CLAUSE_MAP_KIND (nc2)
+				== GOMP_MAP_GRID_STRIDE))
+			  {
+			    stride = OMP_CLAUSE_DECL (nc2);
+			    stride = fold_convert (sizetype, stride);
+			    nc = nc2;
+			  }
+
+			if (tdim >= 0)
+			  {
+			    /* We have an array shape -- use that to find the
+			       total size of the data on the target to look up
+			       in libgomp.  */
+			    tree dtype = TYPE_DOMAIN (tdims[tdim]);
+			    tree minval = TYPE_MIN_VALUE (dtype);
+			    tree maxval = TYPE_MAX_VALUE (dtype);
+			    minval = fold_convert (sizetype, minval);
+			    maxval = fold_convert (sizetype, maxval);
+			    dim = size_binop (MINUS_EXPR, maxval, minval);
+			    dim = size_binop (PLUS_EXPR, dim,
+					      size_one_node);
+			    arrsize = size_binop (MULT_EXPR, arrsize, dim);
+			  }
+			else if (pointer_based && !handled_pointer_section)
+			  {
+			    /* Use the selected array section to determine the
+			       size of the array.  */
+			    tree tmp = size_binop (MULT_EXPR, len, stride);
+			    tmp = size_binop (MINUS_EXPR, tmp, stride);
+			    tmp = size_binop (PLUS_EXPR, tmp, size_one_node);
+			    dim = size_binop (PLUS_EXPR, index, tmp);
+			    arrsize = size_binop (MULT_EXPR, arrsize, dim);
+			    handled_pointer_section = true;
+			  }
+			else
+			  {
+			    if (pointer_based)
+			      error_at (OMP_CLAUSE_LOCATION (c),
+					"too many array section specifiers "
+					"for pointer-based array");
+			    else
+			      error_at (OMP_CLAUSE_LOCATION (c),
+					"too many array section specifiers "
+					"for array");
+			    dim = index = len = stride = error_mark_node;
+			  }
+			tdim--;
+
+			c = nc;
+		      }
+		    else
+		      {
+			/* We have more array dimensions than array section
+			   specifiers.  Copy the whole span.  */
+			tree dtype = TYPE_DOMAIN (tdims[tdim]);
+			tree minval = TYPE_MIN_VALUE (dtype);
+			tree maxval = TYPE_MAX_VALUE (dtype);
+			minval = fold_convert (sizetype, minval);
+			maxval = fold_convert (sizetype, maxval);
+			dim = size_binop (MINUS_EXPR, maxval, minval);
+			dim = size_binop (PLUS_EXPR, dim, size_one_node);
+			len = dim;
+			index = size_zero_node;
+		      }
+
+		    if (TREE_CODE (dim) != INTEGER_CST)
+		      TREE_STATIC (dim_tmp) = 0;
+
+		    if (TREE_CODE (index) != INTEGER_CST)
+		      TREE_STATIC (index_tmp) = 0;
+
+		    if (TREE_CODE (len) != INTEGER_CST)
+		      TREE_STATIC (len_tmp) = 0;
+
+		    if (TREE_CODE (stride) != INTEGER_CST)
+		      TREE_STATIC (stride_tmp) = 0;
+
+		    tree cidx = size_int (i);
+		    CONSTRUCTOR_APPEND_ELT (vdim, cidx, dim);
+		    CONSTRUCTOR_APPEND_ELT (vindex, cidx, index);
+		    CONSTRUCTOR_APPEND_ELT (vlen, cidx, len);
+		    CONSTRUCTOR_APPEND_ELT (vstride, cidx, stride);
+		  }
+
+		/* The size of the whole array -- to make sure we find any
+		   part of the array via splay-tree lookup that might be
+		   mapped on the target at runtime.  */
+		OMP_CLAUSE_SIZE (oc) = arrsize;
+
+		tree cdim = build_constructor (size_arr_type, vdim);
+		tree cindex = build_constructor (size_arr_type, vindex);
+		tree clen = build_constructor (size_arr_type, vlen);
+		tree cstride = build_constructor (size_arr_type, vstride);
+
+		if (TREE_STATIC (dim_tmp))
+		  DECL_INITIAL (dim_tmp) = cdim;
+		else
+		  gimplify_assign (dim_tmp, cdim, &ilist);
+
+		if (TREE_STATIC (index_tmp))
+		  DECL_INITIAL (index_tmp) = cindex;
+		else
+		  gimplify_assign (index_tmp, cindex, &ilist);
+
+		if (TREE_STATIC (len_tmp))
+		  DECL_INITIAL (len_tmp) = clen;
+		else
+		  gimplify_assign (len_tmp, clen, &ilist);
+
+		if (TREE_STATIC (stride_tmp))
+		  DECL_INITIAL (stride_tmp) = cstride;
+		else
+		  gimplify_assign (stride_tmp, cstride, &ilist);
+
+		tree desc_type = TREE_TYPE (desc);
+
+		tree ndims_field = TYPE_FIELDS (desc_type);
+		tree elemsize_field = DECL_CHAIN (ndims_field);
+		tree dim_field = DECL_CHAIN (elemsize_field);
+		tree index_field = DECL_CHAIN (dim_field);
+		tree len_field = DECL_CHAIN (index_field);
+		tree stride_field = DECL_CHAIN (len_field);
+
+		vec<constructor_elt, va_gc> *v;
+		vec_alloc (v, 6);
+
+		bool all_static = (TREE_STATIC (dim_tmp)
+				   && TREE_STATIC (index_tmp)
+				   && TREE_STATIC (len_tmp)
+				   && TREE_STATIC (stride_tmp));
+
+		dim_tmp = build4 (ARRAY_REF, sizetype, dim_tmp, size_zero_node,
+				  NULL_TREE, NULL_TREE);
+		dim_tmp = build_fold_addr_expr (dim_tmp);
+
+		/* TODO: we could skip all-zeros index.  */
+		index_tmp = build4 (ARRAY_REF, sizetype, index_tmp,
+				    size_zero_node, NULL_TREE, NULL_TREE);
+		index_tmp = build_fold_addr_expr (index_tmp);
+
+		len_tmp = build4 (ARRAY_REF, sizetype, len_tmp, size_zero_node,
+				  NULL_TREE, NULL_TREE);
+		len_tmp = build_fold_addr_expr (len_tmp);
+
+		/* TODO: we could skip all-ones stride.  */
+		stride_tmp = build4 (ARRAY_REF, sizetype, stride_tmp,
+				     size_zero_node, NULL_TREE, NULL_TREE);
+		stride_tmp = build_fold_addr_expr (stride_tmp);
+
+		elsize = fold_convert (sizetype, elsize);
+		tree ndims = size_int (dims);
+
+		CONSTRUCTOR_APPEND_ELT (v, ndims_field, ndims);
+		CONSTRUCTOR_APPEND_ELT (v, elemsize_field, elsize);
+		CONSTRUCTOR_APPEND_ELT (v, dim_field, dim_tmp);
+		CONSTRUCTOR_APPEND_ELT (v, index_field, index_tmp);
+		CONSTRUCTOR_APPEND_ELT (v, len_field, len_tmp);
+		CONSTRUCTOR_APPEND_ELT (v, stride_field, stride_tmp);
+
+		tree desc_ctor = build_constructor (desc_type, v);
+
+		if (all_static)
+		  {
+		    TREE_STATIC (desc) = 1;
+		    DECL_INITIAL (desc) = desc_ctor;
+		  }
+		else
+		  gimplify_assign (desc, desc_ctor, &ilist);
+
+		OMP_CLAUSE_CHAIN (dn) = OMP_CLAUSE_CHAIN (nc);
+		c = oc;
+		nc = c;
+	      }
+	    else if (!DECL_P (ovar))
 	      {
 		if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 		    && OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c))
diff --git a/gcc/testsuite/g++.dg/gomp/array-shaping-1.C b/gcc/testsuite/g++.dg/gomp/array-shaping-1.C
new file mode 100644
index 000000000000..8627aa7ffb35
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/array-shaping-1.C
@@ -0,0 +1,22 @@
+// { dg-do compile }
+// { dg-additional-options "-fdump-tree-original" }
+
+template<typename T, typename E, int A, int B, int C, int D>
+void foo ()
+{
+  T *ptr;
+  E a = A, b = B, c = C, d = D;
+
+  /* Dependent types for indices.  */
+#pragma omp target update from(([a][b+1][c][d]) ptr[1:a-2][1:b][1:c-2][1:d-2])
+// { dg-final { scan-tree-dump {map\(from_grid:VIEW_CONVERT_EXPR.*\(\*ptr\) \[len: 1\]\) map\(grid_dim:1 \[len: [^\]]+\]\) map\(grid_dim:1 \[len: [^\]]+\]\) map\(grid_dim:1 \[len: [^\]]+\]\) map\(grid_dim:1 \[len: [^]]+\]\)} "original" } }
+}
+
+int main()
+{
+  char *ptr;
+
+  foo<char, short, 3, 4, 5, 6> ();
+
+  return 0;
+}
diff --git a/gcc/testsuite/g++.dg/gomp/array-shaping-2.C b/gcc/testsuite/g++.dg/gomp/array-shaping-2.C
new file mode 100644
index 000000000000..861d66261a14
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/array-shaping-2.C
@@ -0,0 +1,134 @@
+// { dg-do compile }
+// { dg-additional-options "-fdump-tree-original" }
+
+template<typename T>
+struct St
+{
+  T ***ppptr;
+  T ***&rppptr;
+
+  St(T ***p, T ***&rp) : ppptr(p), rppptr(rp) { }
+};
+
+template<typename A, typename B>
+void foo()
+{
+  A *ptr;
+  A **pptr = &ptr;
+  A ***ppptr = &pptr;
+  A ***&rppptr = ppptr;
+
+#pragma omp target update to(([10]) (**ppptr)[3:4:2])
+// { dg-final { scan-tree-dump {map\(to_grid:VIEW_CONVERT_EXPR<int\[10\]>\(\*\*\*ppptr\) \[len: [0-9]+\]\) map\(grid_dim:3 \[len: 4\]\) map\(grid_stride:2\)} "original" } }
+
+#pragma omp target update to(([10]) (**rppptr)[3:4:2])
+// { dg-final { scan-tree-dump {map\(to_grid:VIEW_CONVERT_EXPR<int\[10\]>\(\*\*\*\*rppptr\) \[len: [0-9]+\]\) map\(grid_dim:3 \[len: 4\]\) map\(grid_stride:2\)} "original" } }
+
+#pragma omp target update to((**ppptr)[3:4:2])
+// { dg-final { scan-tree-dump {map\(to_grid:\*\*ppptr \[len: [0-9]+\]\) map\(grid_dim:3 \[len: 4\]\) map\(grid_stride:2\)} "original" } }
+
+#pragma omp target update to((**rppptr)[3:4:2])
+// { dg-final { scan-tree-dump {map\(to_grid:\*\*\*rppptr \[len: [0-9]+\]\) map\(grid_dim:3 \[len: 4\]\) map\(grid_stride:2\)} "original" } }
+
+  B *ptr2;
+  B **pptr2 = &ptr2;
+  B ***ppptr2 = &pptr2;
+  St<B> *s = new St<B>(ppptr2, ppptr2);
+  St<B> **ps = &s;
+  St<B> **&rps = ps;
+
+#pragma omp target update from(([10]) (**(*ps)->ppptr)[3:4:2])
+// { dg-final { scan-tree-dump {map\(from_grid:VIEW_CONVERT_EXPR<long int\[10\]>\(\*\*\*\(\*ps\)->ppptr\) \[len: [0-9]+\]\) map\(grid_dim:3 \[len: 4\]\) map\(grid_stride:2\)} "original" } }
+
+#pragma omp target update from(([10]) (**(*rps)->rppptr)[3:4:2])
+// { dg-final { scan-tree-dump {map\(from_grid:VIEW_CONVERT_EXPR<long int\[10\]>\(\*\*\*\*\(\*\*rps\)->rppptr\) \[len: [0-9]+\]\) map\(grid_dim:3 \[len: 4\]\) map\(grid_stride:2\)} "original" } }
+
+#pragma omp target update from((**(*ps)->ppptr)[3:4:2])
+// { dg-final { scan-tree-dump {map\(from_grid:\*\*\(\*ps\)->ppptr \[len: [0-9]+\]\) map\(grid_dim:3 \[len: 4\]\) map\(grid_stride:2\)} "original" } }
+
+#pragma omp target update from((**(*rps)->rppptr)[3:4:2])
+// { dg-final { scan-tree-dump {map\(from_grid:\*\*\*\(\*\*rps\)->rppptr \[len: [0-9]+\]\) map\(grid_dim:3 \[len: 4\]\) map\(grid_stride:2\)} "original" } }
+
+  B arr[10][10];
+  B (*parr)[10][10] = &arr;
+  B (**pparr2)[10][10] = &parr;
+  B (**&rpparr2)[10][10] = pparr2;
+
+#pragma omp target update from(**pparr2)
+// { dg-final { scan-tree-dump {from\(\*NON_LVALUE_EXPR <\*pparr2> \[len: [0-9]+\]\)} "original" } }
+
+#pragma omp target update to((**pparr2)[1:5:2][3:4:2])
+// { dg-final { scan-tree-dump {map\(to_grid:\*\*pparr2 \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\) map\(grid_dim:3 \[len: 4\]\) map\(grid_stride:2\)} "original" } }
+
+#pragma omp target update from((**rpparr2)[1:5:2][3:4:2])
+// { dg-final { scan-tree-dump {map\(from_grid:\*\*\*rpparr2 \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\) map\(grid_dim:3 \[len: 4\]\) map\(grid_stride:2\)} "original" } }
+
+  delete s;
+}
+
+struct S
+{
+  short ***ppptr;
+  short ***&rppptr;
+
+  S(short ***p, short ***&rp) : ppptr(p), rppptr(rp) { }
+};
+
+int main()
+{
+  char *ptr;
+  char **pptr = &ptr;
+  char ***ppptr = &pptr;
+  char ***&rppptr = ppptr;
+
+#pragma omp target update to(([10]) (**ppptr)[1:5:2])
+// { dg-final { scan-tree-dump {map\(to_grid:VIEW_CONVERT_EXPR<char\[10\]>\(\*\*\*ppptr\) \[len: 1\]\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\)} "original" } }
+
+#pragma omp target update to(([10]) (**rppptr)[1:5:2])
+// { dg-final { scan-tree-dump {map\(to_grid:VIEW_CONVERT_EXPR<char\[10\]>\(\*\*\*\*rppptr\) \[len: 1\]\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\)} "original" } }
+
+#pragma omp target update to((**ppptr)[1:5:2])
+// { dg-final { scan-tree-dump {map\(to_grid:\*\*ppptr \[len: 1\]\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\)} "original" } }
+
+#pragma omp target update to((**rppptr)[1:5:2])
+// { dg-final { scan-tree-dump {map\(to_grid:\*\*\*rppptr \[len: 1\]\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\)} "original" } }
+
+  short *ptr2;
+  short **pptr2 = &ptr2;
+  short ***ppptr2 = &pptr2;
+  S *s = new S(ppptr2, ppptr2);
+  S **ps = &s;
+  S **&rps = ps;
+
+#pragma omp target update from(([10]) (**(*ps)->ppptr)[1:5:2])
+// { dg-final { scan-tree-dump {map\(from_grid:VIEW_CONVERT_EXPR<short int\[10\]>\(\*\*\*\(\*ps\)->ppptr\) \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\)} "original" } }
+
+#pragma omp target update from(([10]) (**(*rps)->rppptr)[1:5:2])
+// { dg-final { scan-tree-dump {map\(from_grid:VIEW_CONVERT_EXPR<short int\[10\]>\(\*\*\*\*\(\*\*rps\)->rppptr\) \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\)} "original" } }
+
+#pragma omp target update from((**(*ps)->ppptr)[1:5:2])
+// { dg-final { scan-tree-dump {map\(from_grid:\*\*\(\*ps\)->ppptr \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\)} "original" } }
+
+#pragma omp target update from((**(*rps)->rppptr)[1:5:2])
+// { dg-final { scan-tree-dump {map\(from_grid:\*\*\*\(\*\*rps\)->rppptr \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\)} "original" } }
+
+  delete s;
+
+  short arr[10][10];
+  short (*parr)[10][10] = &arr;
+  short (**pparr)[10][10] = &parr;
+  short (**&rpparr)[10][10] = pparr;
+
+#pragma omp target update from(**pparr)
+// { dg-final { scan-tree-dump {from\(\*NON_LVALUE_EXPR <\*pparr> \[len: [0-9]+\]\)} "original" } }
+
+#pragma omp target update to((**pparr)[1:5:2][1:5:2])
+// { dg-final { scan-tree-dump {map\(to_grid:\*\*pparr \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\)} "original" } }
+
+#pragma omp target update from((**rpparr)[1:5:2][1:5:2])
+// { dg-final { scan-tree-dump {map\(from_grid:\*\*\*rpparr \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\)} "original" } }
+
+  foo<int, long> ();
+
+  return 0;
+}
diff --git a/gcc/testsuite/g++.dg/gomp/bad-array-shaping-1.C b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-1.C
new file mode 100644
index 000000000000..1f4e68bc065a
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-1.C
@@ -0,0 +1,47 @@
+// { dg-do compile }
+
+#include <string.h>
+#include <assert.h>
+
+template<typename T, int C, int D>
+void foo (T *w)
+{
+  memset (w, 0, sizeof (T) * 100);
+
+#pragma omp target enter data map(to: w[:100])
+
+  for (int j = 0; j < 10; j++)
+    for (int i = 0; i < 10; i++)
+      w[j * 10 + i] = i + j * 3;
+
+#pragma omp target update to(([C][D]) w[3:2][1:8][0:5])
+// { dg-error "too many array section specifiers for" "" { target *-*-* } .-1 }
+// { dg-error "'#pragma omp target update' must contain at least one 'from' or 'to' clauses" "" { target *-*-* } .-2 }
+
+#pragma omp target exit data map(from: w[:100])
+}
+
+int main()
+{
+  float *arr = new float[100];
+
+  memset (arr, 0, sizeof (float) * 100);
+
+#pragma omp target enter data map(to: arr[:100])
+
+  for (int j = 0; j < 10; j++)
+    for (int i = 0; i < 10; i++)
+      arr[j * 10 + i] = i + j * 3;
+
+#pragma omp target update to(([10][10]) arr[3:2][1:8][0:5])
+// { dg-error "too many array section specifiers for" "" { target *-*-* } .-1 }
+// { dg-error "'#pragma omp target update' must contain at least one 'from' or 'to' clauses" "" { target *-*-* } .-2 }
+
+#pragma omp target exit data map(from: arr[:100])
+
+  foo<float, 5, 20> (arr);
+
+  delete[] arr;
+
+  return 0;
+}
diff --git a/gcc/testsuite/g++.dg/gomp/bad-array-shaping-2.C b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-2.C
new file mode 100644
index 000000000000..d32092925464
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-2.C
@@ -0,0 +1,52 @@
+// { dg-do compile }
+
+#include <string.h>
+#include <assert.h>
+
+template<typename T, int C, int D>
+void foo (T *w)
+{
+  /* This isn't allowed. We get a cascade of errors because it looks a bit
+     like lambda-definition syntax  */
+#pragma omp target enter data map(to: ([C][D]) w[:100])
+  // { dg-error {capture of non-variable 'C'} "" { target *-*-* } .-1 }
+  // { dg-error {expected '\{' before '\[' token} "" { target *-*-* } .-2 }
+  // { dg-warning {lambda expressions only available with} "" { target c++98_only } .-3 }
+  // { dg-error {expected '\)' before 'w'} "" { target *-*-* } .-4 }
+  // { dg-error {does not have pointer or array type} "" { target *-*-* } .-5 }
+
+#pragma omp target exit data map(from: ([C][D]) w[:100])
+  // { dg-error {capture of non-variable 'C'} "" { target *-*-* } .-1 }
+  // { dg-error {expected '\{' before '\[' token} "" { target *-*-* } .-2 }
+  // { dg-warning {lambda expressions only available with} "" { target c++98_only } .-3 }
+  // { dg-error {expected '\)' before 'w'} "" { target *-*-* } .-4 }
+  // { dg-error {does not have pointer or array type} "" { target *-*-* } .-5 }
+}
+
+int main()
+{
+  float *arr = new float[100];
+
+  /* This isn't allowed (as above).  */
+#pragma omp target enter data map(to: ([10][10]) arr[:100])
+  // { dg-error {expected identifier before numeric constant} "" { target *-*-* } .-1 }
+  // { dg-error {expected '\{' before '\[' token} "" { target *-*-* } .-2 }
+  // { dg-warning {lambda expressions only available with} "" { target c++98_only } .-3 }
+  // { dg-error {expected '\)' before 'arr'} "" { target *-*-* } .-4 }
+  // { dg-error {no match for 'operator\[\]' in} "" { target *-*-* } .-5 }
+  // { dg-error {'#pragma omp target enter data' must contain at least one 'map' clause} "" { target *-*-*} .-6 }
+
+#pragma omp target exit data map(from: ([10][10]) arr[:100])
+  // { dg-error {expected identifier before numeric constant} "" { target *-*-* } .-1 }
+  // { dg-error {expected '\{' before '\[' token} "" { target *-*-* } .-2 }
+  // { dg-warning {lambda expressions only available with} "" { target c++98_only } .-3 }
+  // { dg-error {no match for 'operator\[\]' in} "" { target *-*-* } .-4 }
+  // { dg-error {expected '\)' before 'arr'} "" { target *-*-* } .-5 }
+  // { dg-error {'#pragma omp target exit data' must contain at least one 'map' clause} "" { target *-*-* } .-6 }
+
+  foo<float, 5, 20> (arr);
+
+  delete[] arr;
+
+  return 0;
+}
diff --git a/gcc/testsuite/g++.dg/gomp/bad-array-shaping-3.C b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-3.C
new file mode 100644
index 000000000000..90d0a5a80c52
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-3.C
@@ -0,0 +1,53 @@
+// { dg-do compile }
+
+#include <string.h>
+#include <assert.h>
+
+template<typename T>
+void foo (T *w)
+{
+  memset (w, 0, sizeof (T) * 100);
+  int c = 50;
+
+#pragma omp target enter data map(to: w[:100])
+
+  for (int j = 0; j < 10; j++)
+    for (int i = 0; i < 10; i++)
+      w[j * 10 + i] = i + j * 3;
+
+  /* This starts out looking like an array-shape cast.  Make sure it's still
+     parsed as a lambda.  */
+#pragma omp target update to(([c] (T *v) -> T { return v[c]; } (w)))
+  // { dg-message {sorry, unimplemented: unsupported map expression} "" { target *-*-* } .-1 }
+  // { dg-warning {lambda expressions only available with} "" { target c++98_only } .-2 }
+
+#pragma omp target exit data map(from: w[:100])
+}
+
+int main()
+{
+  float *arr = new float[100];
+  int c = 50;
+
+  memset (arr, 0, sizeof (float) * 100);
+
+#pragma omp target enter data map(to: arr[:100])
+
+  for (int j = 0; j < 10; j++)
+    for (int i = 0; i < 10; i++)
+      arr[j * 10 + i] = i + j * 3;
+
+  /* As above.  */
+#pragma omp target update to(([c] (float *v) -> float { return v[c]; } (arr)))
+  // { dg-message {sorry, unimplemented: unsupported map expression} "" { target *-*-* } .-1 }
+  // { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-2 }
+  // { dg-warning {lambda expressions only available with} "" { target c++98_only } .-3 }
+
+#pragma omp target exit data map(from: arr[:100])
+
+  foo<float> (arr);
+
+  delete[] arr;
+
+  return 0;
+}
diff --git a/gcc/testsuite/g++.dg/gomp/bad-array-shaping-4.C b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-4.C
new file mode 100644
index 000000000000..4518f03e9a0c
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-4.C
@@ -0,0 +1,60 @@
+// { dg-do compile }
+
+#include <string.h>
+#include <assert.h>
+
+template<typename T>
+extern T* baz(T*);
+
+template<typename T>
+void foo (T *w)
+{
+  memset (w, 0, sizeof (T) * 100);
+  int c = 50;
+
+#pragma omp target enter data map(to: w[:100])
+
+  for (int j = 0; j < 10; j++)
+    for (int i = 0; i < 10; i++)
+      w[j * 10 + i] = i + j * 3;
+
+  /* No array-shaping inside a function call.  */
+#pragma omp target update to(baz(([10][10]) w))
+  // { dg-error {expected identifier before numeric constant} "" { target *-*-* } .-1 }
+  // { dg-error {expected '\{' before '\[' token} "" { target *-*-* } .-2 }
+  // { dg-warning {lambda expressions only available with} "" { target c++98_only } .-3 }
+  // { dg-error {expected '\)' before 'w'} "" { target *-*-* } .-4 }
+  // { dg-error {no match for 'operator\[\]' in} "" { target *-*-* } .-5 }
+
+#pragma omp target exit data map(from: w[:100])
+}
+
+int main()
+{
+  float *arr = new float[100];
+  int c = 50;
+
+  memset (arr, 0, sizeof (float) * 100);
+
+#pragma omp target enter data map(to: arr[:100])
+
+  for (int j = 0; j < 10; j++)
+    for (int i = 0; i < 10; i++)
+      arr[j * 10 + i] = i + j * 3;
+
+  /* As above.  */
+#pragma omp target update to(baz(([10][10]) arr))
+  // { dg-error {expected identifier before numeric constant} "" { target *-*-* } .-1 }
+  // { dg-error {expected '\{' before '\[' token} "" { target *-*-* } .-2 }
+  // { dg-warning {lambda expressions only available with} "" { target c++98_only } .-3 }
+  // { dg-error {no match for 'operator\[\]' in} "" { target *-*-* } .-4 }
+  // { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-5 }
+
+#pragma omp target exit data map(from: arr[:100])
+
+  foo<float> (arr);
+
+  delete[] arr;
+
+  return 0;
+}
diff --git a/gcc/testsuite/g++.dg/gomp/bad-array-shaping-5.C b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-5.C
new file mode 100644
index 000000000000..25edb9d1d9d3
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-5.C
@@ -0,0 +1,55 @@
+// { dg-do compile }
+// { dg-additional-options "-std=c++14" }
+
+#include <string.h>
+#include <assert.h>
+
+template<typename T>
+void foo (T *w)
+{
+  memset (w, 0, sizeof (T) * 100);
+  int c = 50;
+
+#pragma omp target enter data map(to: w[:100])
+
+  for (int j = 0; j < 10; j++)
+    for (int i = 0; i < 10; i++)
+      w[j * 10 + i] = i + j * 3;
+
+  /* No array-shaping inside a lambda body.  */
+#pragma omp target update to([&](const int d) -> auto& { return ([d][d]) w; } (10))
+// { dg-error {expected '\{' before '\[' token} "" { target *-*-* } .-1 }
+// { dg-error {expected ';' before 'w'} "" { target *-*-* } .-2 }
+// { dg-error {no match for 'operator\[\]' in} "" { target *-*-* } .-3 }
+
+#pragma omp target exit data map(from: w[:100])
+}
+
+int main()
+{
+  float *arr = new float[100];
+  int c = 50;
+
+  memset (arr, 0, sizeof (float) * 100);
+
+#pragma omp target enter data map(to: arr[:100])
+
+  for (int j = 0; j < 10; j++)
+    for (int i = 0; i < 10; i++)
+      arr[j * 10 + i] = i + j * 3;
+
+  /* As above.  */
+#pragma omp target update to([&](const int d) -> auto& { return ([d][d]) arr; } (10))
+// { dg-error {expected '\{' before '\[' token} "" { target *-*-* } .-1 }
+// { dg-error {no match for 'operator\[\]' in} "" { target *-*-* } .-2 }
+// { dg-error {expected ';' before 'arr'} "" { target *-*-* } .-3 }
+// { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-4 }
+
+#pragma omp target exit data map(from: arr[:100])
+
+  foo<float> (arr);
+
+  delete[] arr;
+
+  return 0;
+}
diff --git a/gcc/testsuite/g++.dg/gomp/bad-array-shaping-6.C b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-6.C
new file mode 100644
index 000000000000..e796eaa39a3d
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-6.C
@@ -0,0 +1,59 @@
+// { dg-do compile }
+
+#include <string.h>
+#include <assert.h>
+
+template<typename T>
+void foo (T *w)
+{
+  memset (w, 0, sizeof (T) * 100);
+
+#pragma omp target enter data map(to: w[:100])
+
+  for (int j = 0; j < 10; j++)
+    for (int i = 0; i < 10; i++)
+      w[j * 10 + i] = i + j * 3;
+
+  /* No array-shaping inside a statement expression.  */
+#pragma omp target update to( ({ int d = 10; ([d][d]) w; )} )
+// { dg-error {expected '\{' before '\[' token} "" { target *-*-* } .-1 }
+// { dg-warning {lambda expressions only available with} "" { target c++98_only } .-2 }
+// { dg-error {no match for 'operator\[\]'} "" { target *-*-* } .-3 }
+// { dg-error {expected ';' before 'w'} "" { target *-*-* } .-4 }
+// { dg-error {expected primary-expression before '\)' token} "" { target *-*-* } .-5 }
+// { dg-error {expected '\)' before end of line} "" { target *-*-* } .-6 }
+// { dg-message {sorry, unimplemented: unsupported map expression} "" { target *-*-* } .-7 }
+
+#pragma omp target exit data map(from: w[:100])
+}
+
+int main()
+{
+  float *arr = new float[100];
+
+  memset (arr, 0, sizeof (float) * 100);
+
+#pragma omp target enter data map(to: arr[:100])
+
+  for (int j = 0; j < 10; j++)
+    for (int i = 0; i < 10; i++)
+      arr[j * 10 + i] = i + j * 3;
+
+  /* As above.  */
+#pragma omp target update to( ({ int d = 10; ([d][d]) arr; )} )
+// { dg-error {expected '\{' before '\[' token} "" { target *-*-* } .-1 }
+// { dg-warning {lambda expressions only available with} "" { target c++98_only } .-2 }
+// { dg-error {no match for 'operator\[\]'} "" { target *-*-* } .-3 }
+// { dg-error {expected primary-expression before '\)' token} "" { target *-*-* } .-4 }
+// { dg-error {expected '\)' before end of line} "" { target *-*-* } .-5 }
+// { dg-message {sorry, unimplemented: unsupported map expression} "" { target *-*-* } .-6 }
+// { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-7 }
+
+#pragma omp target exit data map(from: arr[:100])
+
+  foo<float> (arr);
+
+  delete[] arr;
+
+  return 0;
+}
diff --git a/gcc/testsuite/g++.dg/gomp/bad-array-shaping-7.C b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-7.C
new file mode 100644
index 000000000000..362c088c085b
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-7.C
@@ -0,0 +1,48 @@
+// { dg-do compile }
+// { dg-additional-options "-std=c++11" }
+
+#include <new>
+
+template<typename T>
+struct St {
+  T *pp;
+};
+
+template<typename T>
+void foo (T *w)
+{
+  alignas (St<T>) unsigned char buf[sizeof (St<T>)];
+  T *sub1;
+
+  /* No array shaping op in brace initialiser (nonsensical anyway, but make
+     sure it doesn't parse).  */
+#pragma omp target update to( new (buf) St<T> { ([10][10]) sub1 } )
+// { dg-error {expected identifier before numeric constant} "" { target *-*-* } .-1 }
+// { dg-error {expected '\{' before '\[' token} "" { target *-*-* } .-2 }
+// { dg-error {expected '\}' before 'sub1'} "" { target *-*-* } .-3 }
+// { dg-error {expected '\)' before 'sub1'} "" { target *-*-* } .-4 }
+// { dg-error {expected '#pragma omp' clause before '\}' token} "" { target *-*-* } .-5 }
+}
+
+struct S {
+  int *pp;
+};
+
+int main()
+{
+  alignas (S) unsigned char buf[sizeof (S)];
+  int *sub1;
+
+  // As above.
+#pragma omp target update to( new (buf) S { ([10][10]) sub1 } )
+// { dg-error {expected identifier before numeric constant} "" { target *-*-* } .-1 }
+// { dg-error {expected '\{' before '\[' token} "" { target *-*-* } .-2 }
+// { dg-error {expected '\}' before 'sub1'} "" { target *-*-* } .-3 }
+// { dg-error {expected '\)' before 'sub1'} "" { target *-*-* } .-4 }
+// { dg-error {expected '#pragma omp' clause before '\}' token} "" { target *-*-* } .-5 }
+// { dg-error {no match for 'operator\[\]'} "" { target *-*-* } .-6 }
+// { dg-error {could not convert} "" { target *-*-* } .-7 }
+// { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-8 }
+
+  return 0;
+}
diff --git a/gcc/testsuite/g++.dg/gomp/bad-array-shaping-8.C b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-8.C
new file mode 100644
index 000000000000..02d7de6088e0
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-8.C
@@ -0,0 +1,50 @@
+// { dg-do compile }
+
+template<typename T>
+void foo ()
+{
+  T *ptr;
+
+#pragma omp target update to(([5][6][7]) ptr[0:4][0:7][0:7])
+// { dg-error {length '7' with stride '1' above array section size in 'to' clause} "" { target *-*-* } .-1 }
+
+#pragma omp target update to(([5][6][7]) ptr[1:5][0:6][0:7])
+// { dg-error {high bound '6' above array section size in 'to' clause} "" { target *-*-* } .-1 }
+
+  // This one's OK...
+#pragma omp target update from(([100]) ptr[3:33:3])
+
+  // But this is one element out of bounds.
+#pragma omp target update from(([100]) ptr[4:33:3])
+// { dg-error {high bound '101' above array section size in 'from' clause} "" { target *-*-* } .-1 }
+
+#pragma omp target update to(([10][10]) ptr[0:9:-1][0:9])
+// { dg-error {length '9' with stride '-1' above array section size in 'to' clause} "" { target *-*-* } .-1 }
+}
+
+int main()
+{
+  char *ptr;
+
+#pragma omp target update to(([5][6][7]) ptr[0:4][0:7][0:7])
+// { dg-error {length '7' with stride '1' above array section size in 'to' clause} "" { target *-*-* } .-1 }
+// { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-2 }
+
+#pragma omp target update to(([5][6][7]) ptr[1:5][0:6][0:7])
+// { dg-error {high bound '6' above array section size in 'to' clause} "" { target *-*-* } .-1 }
+// { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-2 }
+
+#pragma omp target update from(([100]) ptr[3:33:3])
+
+#pragma omp target update from(([100]) ptr[4:33:3])
+// { dg-error {high bound '101' above array section size in 'from' clause} "" { target *-*-* } .-1 }
+// { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-2 }
+
+#pragma omp target update to(([10][10]) ptr[0:9:-1][0:9])
+// { dg-error {length '9' with stride '-1' above array section size in 'to' clause} "" { target *-*-* } .-1 }
+// { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-2 }
+
+  foo<char> ();
+
+  return 0;
+}
diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc
index 2c2c97ba9515..a80b14349d6a 100644
--- a/gcc/tree-pretty-print.cc
+++ b/gcc/tree-pretty-print.cc
@@ -994,6 +994,18 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
 	case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
 	  pp_string (pp, "attach_zero_length_array_section");
 	  break;
+	case GOMP_MAP_TO_GRID:
+	  pp_string (pp, "to_grid");
+	  break;
+	case GOMP_MAP_FROM_GRID:
+	  pp_string (pp, "from_grid");
+	  break;
+	case GOMP_MAP_GRID_DIM:
+	  pp_string (pp, "grid_dim");
+	  break;
+	case GOMP_MAP_GRID_STRIDE:
+	  pp_string (pp, "grid_stride");
+	  break;
 	case GOMP_MAP_UNSET:
 	  pp_string (pp, "unset");
 	  break;
@@ -2589,6 +2601,11 @@ dump_generic_node (pretty_printer *pp, tree node, int spc, dump_flags_t flags,
       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);
+      if (TREE_OPERAND (node, 3))
+	{
+	  pp_colon (pp);
+	  dump_generic_node (pp,  TREE_OPERAND (node, 3), spc, flags, false);
+	}
       pp_right_bracket (pp);
       break;
 
diff --git a/gcc/tree.def b/gcc/tree.def
index 7ab75b5dc07c..441fcb2a34a0 100644
--- a/gcc/tree.def
+++ b/gcc/tree.def
@@ -1321,7 +1321,7 @@ DEFTREECODE (OMP_ATOMIC_CAPTURE_NEW, "omp_atomic_capture_new", tcc_statement, 2)
 DEFTREECODE (OMP_CLAUSE, "omp_clause", tcc_exceptional, 0)
 
 /* An OpenMP array section.  */
-DEFTREECODE (OMP_ARRAY_SECTION, "omp_array_section", tcc_expression, 3)
+DEFTREECODE (OMP_ARRAY_SECTION, "omp_array_section", tcc_expression, 4)
 
 /* TRANSACTION_EXPR tree code.
    Operand 0: BODY: contains body of the transaction.  */
diff --git a/include/gomp-constants.h b/include/gomp-constants.h
index dd6bfdeffc4e..bf7bab558eaf 100644
--- a/include/gomp-constants.h
+++ b/include/gomp-constants.h
@@ -180,6 +180,9 @@ enum gomp_map_kind
     GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION
       =					(GOMP_MAP_DEEP_COPY | 2),
 
+    GOMP_MAP_TO_GRID =			(GOMP_MAP_DEEP_COPY | 4),
+    GOMP_MAP_FROM_GRID =		(GOMP_MAP_DEEP_COPY | 5),
+
     /* Internal to GCC, not used in libgomp.  */
     /* Do not map, but pointer assign a pointer instead.  */
     GOMP_MAP_FIRSTPRIVATE_POINTER =	(GOMP_MAP_LAST | 1),
@@ -197,7 +200,9 @@ enum gomp_map_kind
     GOMP_MAP_POP_MAPPER_NAME =		(GOMP_MAP_LAST | 6),
     /* Used to hold a TREE_LIST of grouped nodes in an 'omp declare mapper'
        definition (only for Fortran at present).  */
-    GOMP_MAP_MAPPING_GROUP =		(GOMP_MAP_LAST | 7)
+    GOMP_MAP_MAPPING_GROUP =		(GOMP_MAP_LAST | 7),
+    GOMP_MAP_GRID_DIM =			(GOMP_MAP_LAST | 8),
+    GOMP_MAP_GRID_STRIDE =		(GOMP_MAP_LAST | 9),
   };
 
 #define GOMP_MAP_COPY_TO_P(X) \
diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index 8fc9379d1b3c..e0352860fc02 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -1303,6 +1303,20 @@ struct target_mem_desc {
 };
 
 
+/* A rectangular section of an array, for noncontiguous target update
+   operations.  Must be kept in sync with
+   omp-low.cc:omp_noncontig_descriptor_type.  */
+
+typedef struct {
+  size_t ndims;
+  size_t elemsize;
+  size_t *dim;
+  size_t *index;
+  size_t *length;
+  size_t *stride;
+} omp_noncontig_array_desc;
+
+
 typedef struct acc_dispatch_t
 {
   /* Execute.  */
diff --git a/libgomp/target.c b/libgomp/target.c
index f096322a4d27..7a84b50bedbb 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -2103,6 +2103,13 @@ goacc_unmap_vars (struct target_mem_desc *tgt, bool do_copyfrom,
   gomp_unmap_vars_internal (tgt, do_copyfrom, NULL, aq);
 }
 
+static int omp_target_memcpy_rect_worker (void *, const void *, size_t, int,
+					  const size_t *, const size_t *,
+					  const size_t *, const size_t *,
+					  const size_t *, const size_t *,
+					  struct gomp_device_descr *,
+					  struct gomp_device_descr *);
+
 static void
 gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs,
 	     size_t *sizes, void *kinds, bool short_mapkind)
@@ -2125,67 +2132,106 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs,
     }
 
   for (i = 0; i < mapnum; i++)
-    if (sizes[i])
-      {
-	cur_node.host_start = (uintptr_t) hostaddrs[i];
-	cur_node.host_end = cur_node.host_start + sizes[i];
-	splay_tree_key n = splay_tree_lookup (&devicep->mem_map, &cur_node);
-	if (n)
-	  {
-	    int kind = get_kind (short_mapkind, kinds, i);
-	    if (n->host_start > cur_node.host_start
-		|| n->host_end < cur_node.host_end)
-	      {
-		gomp_mutex_unlock (&devicep->lock);
-		gomp_fatal ("Trying to update [%p..%p) object when "
-			    "only [%p..%p) is mapped",
-			    (void *) cur_node.host_start,
-			    (void *) cur_node.host_end,
-			    (void *) n->host_start,
-			    (void *) n->host_end);
-	      }
+    {
+      int kind = get_kind (short_mapkind, kinds, i);
+      if ((kind & typemask) == GOMP_MAP_TO_GRID
+	  || (kind & typemask) == GOMP_MAP_FROM_GRID)
+	{
+	  omp_noncontig_array_desc *desc
+	    = (omp_noncontig_array_desc *) hostaddrs[i + 1];
+	  cur_node.host_start = (uintptr_t) hostaddrs[i];
+	  cur_node.host_end = cur_node.host_start + sizes[i];
+	  assert (sizes[i + 1] == sizeof (omp_noncontig_array_desc));
+	  splay_tree_key n = splay_tree_lookup (&devicep->mem_map, &cur_node);
+	  if (n)
+	    {
+	      if (n->aux && n->aux->attach_count)
+		{
+		  gomp_mutex_unlock (&devicep->lock);
+		  gomp_error ("noncontiguous update with attached pointers");
+		  return;
+		}
+	      void *devaddr = (void *) (n->tgt->tgt_start + n->tgt_offset
+					+ cur_node.host_start
+					- n->host_start);
+	      if ((kind & typemask) == GOMP_MAP_TO_GRID)
+		omp_target_memcpy_rect_worker (devaddr, hostaddrs[i],
+					       desc->elemsize, desc->ndims,
+					       desc->length, desc->stride,
+					       desc->index, desc->index,
+					       desc->dim, desc->dim, devicep,
+					       NULL);
+	      else
+		omp_target_memcpy_rect_worker (hostaddrs[i], devaddr,
+					       desc->elemsize, desc->ndims,
+					       desc->length, desc->stride,
+					       desc->index, desc->index,
+					       desc->dim, desc->dim, NULL,
+					       devicep);
+	    }
+	  i++;
+	}
+      else if (sizes[i])
+	{
+	  cur_node.host_start = (uintptr_t) hostaddrs[i];
+	  cur_node.host_end = cur_node.host_start + sizes[i];
+	  splay_tree_key n = splay_tree_lookup (&devicep->mem_map, &cur_node);
+	  if (n)
+	    {
+	      if (n->host_start > cur_node.host_start
+		  || n->host_end < cur_node.host_end)
+		{
+		  gomp_mutex_unlock (&devicep->lock);
+		  gomp_fatal ("Trying to update [%p..%p) object when "
+			      "only [%p..%p) is mapped",
+			      (void *) cur_node.host_start,
+			      (void *) cur_node.host_end,
+			      (void *) n->host_start,
+			      (void *) n->host_end);
+		}
 
-	    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 (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);
-	      }
-	  }
-      }
+		  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);
 }
 
@@ -4524,6 +4570,7 @@ omp_target_memcpy_async (void *dst, const void *src, size_t length,
 static int
 omp_target_memcpy_rect_worker (void *dst, const void *src, size_t element_size,
 			       int num_dims, const size_t *volume,
+			       const size_t *strides,
 			       const size_t *dst_offsets,
 			       const size_t *src_offsets,
 			       const size_t *dst_dimensions,
@@ -4536,7 +4583,7 @@ omp_target_memcpy_rect_worker (void *dst, const void *src, size_t element_size,
   size_t j, dst_off, src_off, length;
   int i, ret;
 
-  if (num_dims == 1)
+  if (num_dims == 1 && (!strides || strides[0] == 1))
     {
       if (__builtin_mul_overflow (element_size, volume[0], &length)
 	  || __builtin_mul_overflow (element_size, dst_offsets[0], &dst_off)
@@ -4567,6 +4614,38 @@ omp_target_memcpy_rect_worker (void *dst, const void *src, size_t element_size,
 	ret = 0;
       return ret ? 0 : EINVAL;
     }
+  else if (num_dims == 1 && strides)
+    {
+      size_t stride;
+
+      assert ((src_devicep == NULL || dst_devicep == NULL)
+	      && (src_devicep != NULL || dst_devicep != NULL));
+
+      if (__builtin_mul_overflow (element_size, dst_offsets[0], &dst_off)
+	  || __builtin_mul_overflow (element_size, src_offsets[0], &src_off))
+	return EINVAL;
+
+      if (strides
+	  && __builtin_mul_overflow (element_size, strides[0], &stride))
+	return EINVAL;
+
+      for (i = 0, ret = 1; i < volume[0] && ret; i++)
+	{
+	  if (src_devicep == NULL)
+	    ret = dst_devicep->host2dev_func (dst_devicep->target_id,
+					      (char *) dst + dst_off,
+					      (const char *) src + src_off,
+					      element_size);
+	  else if (dst_devicep == NULL)
+	    ret = src_devicep->dev2host_func (src_devicep->target_id,
+					      (char *) dst + dst_off,
+					      (const char *) src + src_off,
+					      element_size);
+	  dst_off += stride;
+	  src_off += stride;
+	}
+      return ret ? 0 : EINVAL;
+    }
 
   /* FIXME: it would be nice to have some plugin function to handle
      num_dims == 2 and num_dims == 3 more efficiently.  Larger ones can
@@ -4580,13 +4659,19 @@ omp_target_memcpy_rect_worker (void *dst, const void *src, size_t element_size,
   if (__builtin_mul_overflow (dst_slice, dst_offsets[0], &dst_off)
       || __builtin_mul_overflow (src_slice, src_offsets[0], &src_off))
     return EINVAL;
+  if (strides
+      && (__builtin_mul_overflow (dst_slice, strides[0], &dst_slice)
+	  || __builtin_mul_overflow (src_slice, strides[0], &src_slice)))
+    return EINVAL;
   for (j = 0; j < volume[0]; j++)
     {
       ret = omp_target_memcpy_rect_worker ((char *) dst + dst_off,
 					   (const char *) src + src_off,
 					   element_size, num_dims - 1,
-					   volume + 1, dst_offsets + 1,
-					   src_offsets + 1, dst_dimensions + 1,
+					   volume + 1,
+					   strides ? strides + 1 : NULL,
+					   dst_offsets + 1, src_offsets + 1,
+					   dst_dimensions + 1,
 					   src_dimensions + 1, dst_devicep,
 					   src_devicep);
       if (ret)
@@ -4632,9 +4717,10 @@ omp_target_memcpy_rect_copy (void *dst, const void *src,
   else if (dst_devicep)
     gomp_mutex_lock (&dst_devicep->lock);
   int ret = omp_target_memcpy_rect_worker (dst, src, element_size, num_dims,
-					   volume, dst_offsets, src_offsets,
-					   dst_dimensions, src_dimensions,
-					   dst_devicep, src_devicep);
+					   volume, NULL, dst_offsets,
+					   src_offsets, dst_dimensions,
+					   src_dimensions, dst_devicep,
+					   src_devicep);
   if (src_devicep)
     gomp_mutex_unlock (&src_devicep->lock);
   else if (dst_devicep)
diff --git a/libgomp/testsuite/libgomp.c++/array-shaping-1.C b/libgomp/testsuite/libgomp.c++/array-shaping-1.C
new file mode 100644
index 000000000000..6ff5f9475f6b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/array-shaping-1.C
@@ -0,0 +1,469 @@
+// { dg-do run { target offload_device_nonshared_as } }
+
+#include <string.h>
+#include <assert.h>
+
+volatile int yy = 4, zz = 2, str_str = 2;
+
+template<typename T>
+void foo()
+{
+  T *arr;
+  int x = 5;
+  T arr2d[10][10];
+
+  arr = new T[100];
+
+  /* Update whole reshaped array.  */
+
+  memset (arr, 0, 100 * sizeof (T));
+
+#pragma omp target enter data map(to: arr[:100])
+
+  for (int j = 0; j < x; j++)
+    for (int i = 0; i < 10; i++)
+      arr[j * 10 + i] = i ^ j;
+
+#pragma omp target update to(([10][x]) arr)
+
+#pragma omp target exit data map(from: arr[:100])
+
+  for (int j = 0; j < 10; j++)
+    for (int i = 0; i < 10; i++)
+      if (j < x)
+	assert (arr[j * 10 + i] == i ^ j);
+      else
+	assert (arr[j * 10 + i] == 0);
+
+
+  /* Strided update.  */
+
+  memset (arr, 0, 100 * sizeof (T));
+
+#pragma omp target enter data map(to: arr[:100])
+
+  for (int j = 0; j < 20; j++)
+    for (int i = 0; i < 5; i++)
+      arr[j * 5 + i] = i + j;
+
+#pragma omp target update to(([5][5]) arr[0:3][0:3:2])
+
+#pragma omp target exit data map(from: arr[:100])
+
+  for (int j = 0; j < 20; j++)
+    for (int i = 0; i < 5; i++)
+      if (j < 3 && (i & 1) == 0 && i < 6)
+	assert (arr[j * 5 + i] == i + j);
+      else
+	assert (arr[j * 5 + i] == 0);
+
+
+  /* Reshaped update, contiguous.  */
+
+  memset (arr, 0, 100 * sizeof (T));
+
+#pragma omp target enter data map(to: arr[:100])
+
+  for (int j = 0; j < 20; j++)
+    for (int i = 0; i < 5; i++)
+      arr[j * 5 + i] = 2 * j + i;
+
+#pragma omp target update to(([5][5]) arr[0:5][0:5])
+
+#pragma omp target exit data map(from: arr[:100])
+
+  for (int j = 0; j < 20; j++)
+    for (int i = 0; i < 5; i++)
+      if (j < 5 && i < 5)
+	assert (arr[j * 5 + i] == 2 * j + i);
+      else
+	assert (arr[j * 5 + i] == 0);
+
+
+  /* Strided update on actual array.  */
+
+  memset (arr2d, 0, 100 * sizeof (T));
+
+#pragma omp target enter data map(to: arr2d)
+
+  for (int j = 0; j < 10; j++)
+    for (int i = 0; i < 10; i++)
+      arr2d[j][i] = j + 2 * i;
+
+#pragma omp target update to(arr2d[0:5:2][5:2])
+
+#pragma omp target exit data map(from: arr2d)
+
+  for (int j = 0; j < 10; j++)
+    for (int i = 0; i < 10; i++)
+      if ((j & 1) == 0 && i >= 5 && i < 7)
+	assert (arr2d[j][i] == j + 2 * i);
+      else
+	assert (arr2d[j][i] == 0);
+
+
+  /* Update with non-constant bounds.  */
+
+  memset (arr, 0, 100 * sizeof (T));
+
+#pragma omp target enter data map(to: arr[:100])
+
+  for (int j = 0; j < 10; j++)
+    for (int i = 0; i < 10; i++)
+      arr[j * 10 + i] = (2 * j) ^ i;
+
+  x = 3;
+  int y = yy, z = zz, str = str_str;
+  /* This is actually [0:3:2] [4:2:2].  */
+#pragma omp target update to(([10][10]) arr[0:x:2][y:z:str])
+
+#pragma omp target exit data map(from: arr[:100])
+
+  for (int j = 0; j < 10; j++)
+    for (int i = 0; i < 10; i++)
+      if ((j & 1) == 0 && j < 6 && (i & 1) == 0 && i >= 4 && i < 8)
+	assert (arr[j * 10 + i] == (2 * j) ^ i);
+      else
+	assert (arr[j * 10 + i] == 0);
+
+
+  /* Update with full "major" dimension.  */
+
+  memset (arr, 0, 100 * sizeof (T));
+
+#pragma omp target enter data map(to: arr[:100])
+
+  for (int j = 0; j < 10; j++)
+    for (int i = 0; i < 10; i++)
+      arr[j * 10 + i] = i + j;
+
+#pragma omp target update to(([10][10]) arr[0:10][3:1])
+
+#pragma omp target exit data map(from: arr[:100])
+
+  for (int j = 0; j < 10; j++)
+    for (int i = 0; i < 10; i++)
+      if (i == 3)
+	assert (arr[j * 10 + i] == i + j);
+      else
+	assert (arr[j * 10 + i] == 0);
+
+
+  /* Update with full "minor" dimension.  */
+
+  memset (arr, 0, 100 * sizeof (T));
+
+#pragma omp target enter data map(to: arr[:100])
+
+  for (int j = 0; j < 10; j++)
+    for (int i = 0; i < 10; i++)
+      arr[j * 10 + i] = 3 * (i + j);
+
+#pragma omp target update to(([10][10]) arr[3:2][0:10])
+
+#pragma omp target exit data map(from: arr[:100])
+
+  for (int j = 0; j < 10; j++)
+    for (int i = 0; i < 10; i++)
+      if (j >= 3 && j < 5)
+	assert (arr[j * 10 + i] == 3 * (i + j));
+      else
+	assert (arr[j * 10 + i] == 0);
+
+
+  /* Rectangle update.  */
+
+  memset (arr, 0, 100 * sizeof (T));
+
+#pragma omp target enter data map(to: arr[:100])
+
+  for (int j = 0; j < 10; j++)
+    for (int i = 0; i < 10; i++)
+      arr[j * 10 + i] = 5 * (i + j);
+
+#pragma omp target update to(([10][10]) arr[3:2][0:9])
+
+#pragma omp target exit data map(from: arr[:100])
+
+  for (int j = 0; j < 10; j++)
+    for (int i = 0; i < 10; i++)
+      if (j >= 3 && j < 5 && i < 9)
+	assert (arr[j * 10 + i] == 5 * (i + j));
+      else
+	assert (arr[j * 10 + i] == 0);
+
+
+  /* One-dimensional strided update.  */
+
+  memset (arr, 0, 100 * sizeof (T));
+
+#pragma omp target enter data map(to: arr[:100])
+
+  for (int i = 0; i < 100; i++)
+    arr[i] = i + 99;
+
+#pragma omp target update to(([100]) arr[3:33:3])
+
+#pragma omp target exit data map(from: arr[:100])
+
+  for (int i = 0; i < 100; i++)
+    if (i >= 3 && ((i - 3) % 3) == 0)
+      assert (arr[i] == i + 99);
+    else
+      assert (arr[i] == 0);
+
+
+  /* One-dimensional strided update without explicit array shape.  */
+
+  memset (arr, 0, 100 * sizeof (T));
+
+#pragma omp target enter data map(to: arr[:100])
+
+  for (int i = 0; i < 100; i++)
+    arr[i] = i + 121;
+
+#pragma omp target update to(arr[3:33:3])
+
+#pragma omp target exit data map(from: arr[:100])
+
+  for (int i = 0; i < 100; i++)
+    if (i >= 3 && ((i - 3) % 3) == 0)
+      assert (arr[i] == i + 121);
+    else
+      assert (arr[i] == 0);
+
+  delete[] arr;
+}
+
+int main()
+{
+  int *arr;
+  int x = 5;
+  int arr2d[10][10];
+
+  arr = new int[100];
+
+  /* Update whole reshaped array.  */
+
+  memset (arr, 0, 100 * sizeof (int));
+
+#pragma omp target enter data map(to: arr[:100])
+
+  for (int j = 0; j < x; j++)
+    for (int i = 0; i < 10; i++)
+      arr[j * 10 + i] = i ^ j;
+
+#pragma omp target update to(([10][x]) arr)
+
+#pragma omp target exit data map(from: arr[:100])
+
+  for (int j = 0; j < 10; j++)
+    for (int i = 0; i < 10; i++)
+      if (j < x)
+	assert (arr[j * 10 + i] == i ^ j);
+      else
+	assert (arr[j * 10 + i] == 0);
+
+
+  /* Strided update.  */
+
+  memset (arr, 0, 100 * sizeof (int));
+
+#pragma omp target enter data map(to: arr[:100])
+
+  for (int j = 0; j < 20; j++)
+    for (int i = 0; i < 5; i++)
+      arr[j * 5 + i] = i + j;
+
+#pragma omp target update to(([5][5]) arr[0:3][0:3:2])
+
+#pragma omp target exit data map(from: arr[:100])
+
+  for (int j = 0; j < 20; j++)
+    for (int i = 0; i < 5; i++)
+      if (j < 3 && (i & 1) == 0 && i < 6)
+	assert (arr[j * 5 + i] == i + j);
+      else
+	assert (arr[j * 5 + i] == 0);
+
+
+  /* Reshaped update, contiguous.  */
+
+  memset (arr, 0, 100 * sizeof (int));
+
+#pragma omp target enter data map(to: arr[:100])
+
+  for (int j = 0; j < 20; j++)
+    for (int i = 0; i < 5; i++)
+      arr[j * 5 + i] = 2 * j + i;
+
+#pragma omp target update to(([5][5]) arr[0:5][0:5])
+
+#pragma omp target exit data map(from: arr[:100])
+
+  for (int j = 0; j < 20; j++)
+    for (int i = 0; i < 5; i++)
+      if (j < 5 && i < 5)
+	assert (arr[j * 5 + i] == 2 * j + i);
+      else
+	assert (arr[j * 5 + i] == 0);
+
+
+  /* Strided update on actual array.  */
+
+  memset (arr2d, 0, 100 * sizeof (int));
+
+#pragma omp target enter data map(to: arr2d)
+
+  for (int j = 0; j < 10; j++)
+    for (int i = 0; i < 10; i++)
+      arr2d[j][i] = j + 2 * i;
+
+#pragma omp target update to(arr2d[0:5:2][5:2])
+
+#pragma omp target exit data map(from: arr2d)
+
+  for (int j = 0; j < 10; j++)
+    for (int i = 0; i < 10; i++)
+      if ((j & 1) == 0 && i >= 5 && i < 7)
+	assert (arr2d[j][i] == j + 2 * i);
+      else
+	assert (arr2d[j][i] == 0);
+
+
+  /* Update with non-constant bounds.  */
+
+  memset (arr, 0, 100 * sizeof (int));
+
+#pragma omp target enter data map(to: arr[:100])
+
+  for (int j = 0; j < 10; j++)
+    for (int i = 0; i < 10; i++)
+      arr[j * 10 + i] = (2 * j) ^ i;
+
+  x = 3;
+  int y = yy, z = zz, str = str_str;
+  /* This is actually [0:3:2] [4:2:2].  */
+#pragma omp target update to(([10][10]) arr[0:x:2][y:z:str])
+
+#pragma omp target exit data map(from: arr[:100])
+
+  for (int j = 0; j < 10; j++)
+    for (int i = 0; i < 10; i++)
+      if ((j & 1) == 0 && j < 6 && (i & 1) == 0 && i >= 4 && i < 8)
+	assert (arr[j * 10 + i] == (2 * j) ^ i);
+      else
+	assert (arr[j * 10 + i] == 0);
+
+
+  /* Update with full "major" dimension.  */
+
+  memset (arr, 0, 100 * sizeof (int));
+
+#pragma omp target enter data map(to: arr[:100])
+
+  for (int j = 0; j < 10; j++)
+    for (int i = 0; i < 10; i++)
+      arr[j * 10 + i] = i + j;
+
+#pragma omp target update to(([10][10]) arr[0:10][3:1])
+
+#pragma omp target exit data map(from: arr[:100])
+
+  for (int j = 0; j < 10; j++)
+    for (int i = 0; i < 10; i++)
+      if (i == 3)
+	assert (arr[j * 10 + i] == i + j);
+      else
+	assert (arr[j * 10 + i] == 0);
+
+
+  /* Update with full "minor" dimension.  */
+
+  memset (arr, 0, 100 * sizeof (int));
+
+#pragma omp target enter data map(to: arr[:100])
+
+  for (int j = 0; j < 10; j++)
+    for (int i = 0; i < 10; i++)
+      arr[j * 10 + i] = 3 * (i + j);
+
+#pragma omp target update to(([10][10]) arr[3:2][0:10])
+
+#pragma omp target exit data map(from: arr[:100])
+
+  for (int j = 0; j < 10; j++)
+    for (int i = 0; i < 10; i++)
+      if (j >= 3 && j < 5)
+	assert (arr[j * 10 + i] == 3 * (i + j));
+      else
+	assert (arr[j * 10 + i] == 0);
+
+
+  /* Rectangle update.  */
+
+  memset (arr, 0, 100 * sizeof (int));
+
+#pragma omp target enter data map(to: arr[:100])
+
+  for (int j = 0; j < 10; j++)
+    for (int i = 0; i < 10; i++)
+      arr[j * 10 + i] = 5 * (i + j);
+
+#pragma omp target update to(([10][10]) arr[3:2][0:9])
+
+#pragma omp target exit data map(from: arr[:100])
+
+  for (int j = 0; j < 10; j++)
+    for (int i = 0; i < 10; i++)
+      if (j >= 3 && j < 5 && i < 9)
+	assert (arr[j * 10 + i] == 5 * (i + j));
+      else
+	assert (arr[j * 10 + i] == 0);
+
+
+  /* One-dimensional strided update.  */
+
+  memset (arr, 0, 100 * sizeof (int));
+
+#pragma omp target enter data map(to: arr[:100])
+
+  for (int i = 0; i < 100; i++)
+    arr[i] = i + 99;
+
+#pragma omp target update to(([100]) arr[3:33:3])
+
+#pragma omp target exit data map(from: arr[:100])
+
+  for (int i = 0; i < 100; i++)
+    if (i >= 3 && ((i - 3) % 3) == 0)
+      assert (arr[i] == i + 99);
+    else
+      assert (arr[i] == 0);
+
+
+  /* One-dimensional strided update without explicit array shape.  */
+
+  memset (arr, 0, 100 * sizeof (int));
+
+#pragma omp target enter data map(to: arr[:100])
+
+  for (int i = 0; i < 100; i++)
+    arr[i] = i + 121;
+
+#pragma omp target update to(arr[3:33:3])
+
+#pragma omp target exit data map(from: arr[:100])
+
+  for (int i = 0; i < 100; i++)
+    if (i >= 3 && ((i - 3) % 3) == 0)
+      assert (arr[i] == i + 121);
+    else
+      assert (arr[i] == 0);
+
+  delete[] arr;
+
+  foo<long> ();
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/array-shaping-10.C b/libgomp/testsuite/libgomp.c++/array-shaping-10.C
new file mode 100644
index 000000000000..648f02d34798
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/array-shaping-10.C
@@ -0,0 +1,61 @@
+// { dg-do run { target offload_device_nonshared_as } }
+
+#include <assert.h>
+#include <string.h>
+
+#define N 10
+
+template<typename T>
+void foo ()
+{
+  T tarr[N * N];
+
+  memset (tarr, 0, N * N * sizeof (T));
+
+#pragma omp target enter data map(to: tarr)
+
+#pragma omp target
+  {
+    for (int i = 0; i < N; i++)
+      for (int j = 0; j < N; j++)
+	tarr[i * N + j] = 2 * (i + j);
+  }
+
+  /* An array, but cast to a pointer, then reshaped.  */
+#pragma omp target update from(([N][N]) ((T *) &tarr[0])[4:3][5:3])
+
+  for (int i = 4; i < 7; i++)
+    for (int j = 5; j < 8; j++)
+      assert (tarr[i * N + j] == 2 * (i + j));
+
+#pragma omp target exit data map(delete: tarr)
+}
+
+int main ()
+{
+  int iarr[N * N];
+
+  memset (iarr, 0, N * N * sizeof (int));
+
+#pragma omp target enter data map(to: iarr)
+
+#pragma omp target
+  {
+    for (int i = 0; i < 10; i++)
+      for (int j = 0; j < 10; j++)
+	iarr[i * 10 + j] = i + j;
+  }
+
+  /* An array, but cast to a pointer, then reshaped.  */
+#pragma omp target update from(([10][10]) ((int *) &iarr[0])[4:3][4:3])
+
+  for (int i = 4; i < 7; i++)
+    for (int j = 4; j < 7; j++)
+      assert (iarr[i * 10 + j] == i + j);
+
+#pragma omp target exit data map(delete: iarr)
+
+  foo<unsigned short> ();
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/array-shaping-11.C b/libgomp/testsuite/libgomp.c++/array-shaping-11.C
new file mode 100644
index 000000000000..6b15bd62fb1f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/array-shaping-11.C
@@ -0,0 +1,63 @@
+// { dg-do run { target offload_device_nonshared_as } }
+
+#include <assert.h>
+#include <string.h>
+
+#define N 10
+
+template<typename T>
+void foo ()
+{
+  T tarr_real[N * N];
+  T (&tarr)[N * N] = tarr_real;
+
+  memset (tarr, 0, N * N * sizeof (T));
+
+#pragma omp target enter data map(to: tarr)
+
+#pragma omp target
+  {
+    for (int i = 0; i < N; i++)
+      for (int j = 0; j < N; j++)
+	tarr[i * N + j] = 2 * (i + j);
+  }
+
+  /* A ref to an array, but cast to a pointer, then reshaped.  */
+#pragma omp target update from(([N][N]) ((T *) &tarr[0])[4:3][5:3])
+
+  for (int i = 4; i < 7; i++)
+    for (int j = 5; j < 8; j++)
+      assert (tarr[i * N + j] == 2 * (i + j));
+
+#pragma omp target exit data map(delete: tarr)
+}
+
+int main ()
+{
+  int iarr_real[N * N];
+  int (&iarr)[N * N] = iarr_real;
+
+  memset (iarr, 0, N * N * sizeof (int));
+
+#pragma omp target enter data map(to: iarr)
+
+#pragma omp target
+  {
+    for (int i = 0; i < 10; i++)
+      for (int j = 0; j < 10; j++)
+	iarr[i * 10 + j] = i + j;
+  }
+
+  /* A ref to an array, but cast to a pointer, then reshaped.  */
+#pragma omp target update from(([10][10]) ((int *) &iarr[0])[4:3][4:3])
+
+  for (int i = 4; i < 7; i++)
+    for (int j = 4; j < 7; j++)
+      assert (iarr[i * 10 + j] == i + j);
+
+#pragma omp target exit data map(delete: iarr)
+
+  foo<unsigned short> ();
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/array-shaping-12.C b/libgomp/testsuite/libgomp.c++/array-shaping-12.C
new file mode 100644
index 000000000000..103c99aa847e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/array-shaping-12.C
@@ -0,0 +1,65 @@
+// { dg-do run { target offload_device_nonshared_as } }
+
+#include <assert.h>
+#include <string.h>
+
+#define N 10
+
+template<typename T>
+void foo ()
+{
+  T tarr_real[N * N];
+  T *tarrp = &tarr_real[0];
+  T **tarrpp = &tarrp;
+
+  memset (tarrp, 0, N * N * sizeof (T));
+
+#pragma omp target enter data map(to: tarr_real)
+
+#pragma omp target
+  {
+    for (int i = 0; i < N; i++)
+      for (int j = 0; j < N; j++)
+	tarrp[i * N + j] = 2 * (i + j);
+  }
+
+  /* A pointer with an extra indirection.  */
+#pragma omp target update from(([N][N]) (*tarrpp)[4:3][5:3])
+
+  for (int i = 4; i < 7; i++)
+    for (int j = 5; j < 8; j++)
+      assert (tarrp[i * N + j] == 2 * (i + j));
+
+#pragma omp target exit data map(delete: tarr_real)
+}
+
+int main ()
+{
+  int iarr_real[N * N];
+  int *iarrp = &iarr_real[0];
+  int **iarrpp = &iarrp;
+
+  memset (iarrp, 0, N * N * sizeof (int));
+
+#pragma omp target enter data map(to: iarr_real)
+
+#pragma omp target
+  {
+    for (int i = 0; i < 10; i++)
+      for (int j = 0; j < 10; j++)
+	iarrp[i * 10 + j] = i + j;
+  }
+
+  /* A pointer with an extra indirection.  */
+#pragma omp target update from(([10][10]) (*iarrpp)[4:3][4:3])
+
+  for (int i = 4; i < 7; i++)
+    for (int j = 4; j < 7; j++)
+      assert (iarrp[i * 10 + j] == i + j);
+
+#pragma omp target exit data map(delete: iarr_real)
+
+  foo<unsigned short> ();
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/array-shaping-13.C b/libgomp/testsuite/libgomp.c++/array-shaping-13.C
new file mode 100644
index 000000000000..29345ca4264c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/array-shaping-13.C
@@ -0,0 +1,89 @@
+// { dg-do run { target offload_device_nonshared_as } }
+
+#include <assert.h>
+#include <string.h>
+
+#define N 10
+
+template<typename T>
+void foo ()
+{
+  T *tptr = new T[N * N * N];
+
+  memset (tptr, 0, N * N * N * sizeof (T));
+
+#pragma omp target enter data map(to: tptr[0:N*N*N])
+
+#pragma omp target
+  {
+    for (int i = 0; i < N; i++)
+      for (int j = 0; j < N; j++)
+	tptr[i * N * N + 4 * N + j] = 2 * (i + j);
+  }
+
+  /* An array ref between two array sections.  */
+#pragma omp target update from(([N][N][N]) tptr[4:3][4][5:3])
+
+  for (int i = 4; i < 7; i++)
+    for (int j = 5; j < 8; j++)
+      assert (tptr[i * N * N + 4 * N + j] == 2 * (i + j));
+
+  memset (tptr, 0, N * N * N * sizeof (T));
+
+  for (int i = 0; i < N; i++)
+    tptr[2 * N * N + i * N + 4] = 4 * i;
+
+  /* Array section between two array refs.  */
+#pragma omp target update to(([N][N][N]) tptr[2][3:6][4])
+
+#pragma omp target exit data map(from: tptr[0:N*N*N])
+
+  for (int i = 3; i < 9; i++)
+    assert (tptr[2 * N * N + i * N + 4] == 4 * i);
+
+#pragma omp target exit data map(delete: tptr[0:N*N*N])
+
+  delete[] tptr;
+}
+
+int main ()
+{
+  int *iptr = new int[N * N * N];
+
+  memset (iptr, 0, N * N * N * sizeof (int));
+
+#pragma omp target enter data map(to: iptr[0:N*N*N])
+
+#pragma omp target
+  {
+    for (int i = 0; i < N; i++)
+      for (int j = 0; j < N; j++)
+	iptr[i * N * N + 4 * N + j] = i + j;
+  }
+
+  /* An array ref between two array sections.  */
+#pragma omp target update from(([N][N][N]) iptr[2:3][4][6:3])
+
+  for (int i = 2; i < 5; i++)
+    for (int j = 6; j < 9; j++)
+      assert (iptr[i * N * N + 4 * N + j] == i + j);
+
+  memset (iptr, 0, N * N * N * sizeof (int));
+
+  for (int i = 0; i < N; i++)
+    iptr[2 * N * N + i * N + 4] = 3 * i;
+
+  /* Array section between two array refs.  */
+#pragma omp target update to(([N][N][N]) iptr[2][3:6][4])
+
+#pragma omp target exit data map(from: iptr[0:N*N*N])
+
+  for (int i = 3; i < 9; i++)
+    assert (iptr[2 * N * N + i * N + 4] == 3 * i);
+
+  delete[] iptr;
+
+  foo<unsigned long> ();
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/array-shaping-2.C b/libgomp/testsuite/libgomp.c++/array-shaping-2.C
new file mode 100644
index 000000000000..027543e8d297
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/array-shaping-2.C
@@ -0,0 +1,38 @@
+// { dg-do run { target offload_device_nonshared_as } }
+
+#include <string.h>
+#include <assert.h>
+
+template<typename T>
+void foo (T *w)
+{
+  memset (w, 0, sizeof (T) * 100);
+
+#pragma omp target enter data map(to: w[:100])
+
+  for (int j = 0; j < 10; j++)
+    for (int i = 0; i < 10; i++)
+      w[j * 10 + i] = i + j;
+
+#pragma omp target update to(([10][10]) w[3:2][1:8])
+
+#pragma omp target exit data map(from: w[:100])
+
+  for (int j = 0; j < 10; j++)
+    for (int i = 0; i < 10; i++)
+      if (j >= 3 && j < 5 && i >= 1 && i < 9)
+	assert (w[j * 10 + i] == i + j);
+      else
+	assert (w[j * 10 + i] == 0);
+}
+
+int main()
+{
+  int *arr = new int[100];
+
+  foo<int> (arr);
+
+  delete[] arr;
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/array-shaping-3.C b/libgomp/testsuite/libgomp.c++/array-shaping-3.C
new file mode 100644
index 000000000000..09ff04bc1145
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/array-shaping-3.C
@@ -0,0 +1,38 @@
+// { dg-do run { target offload_device_nonshared_as } }
+
+#include <string.h>
+#include <assert.h>
+
+template<int C, int D>
+void foo (double *w)
+{
+  memset (w, 0, sizeof (double) * 100);
+
+#pragma omp target enter data map(to: w[:100])
+
+  for (int j = 0; j < 10; j++)
+    for (int i = 0; i < 10; i++)
+      w[j * 10 + i] = i * 3 + j * 2;
+
+#pragma omp target update to(([C][D]) w[3:2][1:8])
+
+#pragma omp target exit data map(from: w[:100])
+
+  for (int j = 0; j < 10; j++)
+    for (int i = 0; i < 10; i++)
+      if (j >= 3 && j < 5 && i >= 1 && i < 9)
+	assert (w[j * 10 + i] == i * 3 + j * 2);
+      else
+	assert (w[j * 10 + i] == 0.0f);
+}
+
+int main()
+{
+  double *arr = new double[100];
+
+  foo<10, 10> (arr);
+
+  delete[] arr;
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/array-shaping-4.C b/libgomp/testsuite/libgomp.c++/array-shaping-4.C
new file mode 100644
index 000000000000..efa115e8be6b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/array-shaping-4.C
@@ -0,0 +1,38 @@
+// { dg-do run { target offload_device_nonshared_as } }
+
+#include <string.h>
+#include <assert.h>
+
+template<auto C, auto D>
+void foo (double *w)
+{
+  memset (w, 0, sizeof (double) * 100);
+
+#pragma omp target enter data map(to: w[:100])
+
+  for (int j = 0; j < 10; j++)
+    for (int i = 0; i < 10; i++)
+      w[j * 10 + i] = i * 2 + j * 3;
+
+#pragma omp target update to(([C][D]) w[3:2][1:8])
+
+#pragma omp target exit data map(from: w[:100])
+
+  for (int j = 0; j < 10; j++)
+    for (int i = 0; i < 10; i++)
+      if (j >= 3 && j < 5 && i >= 1 && i < 9)
+	assert (w[j * 10 + i] == i * 2 + j * 3);
+      else
+	assert (w[j * 10 + i] == 0.0f);
+}
+
+int main()
+{
+  double *arr = new double[100];
+
+  foo<10, 10> (arr);
+
+  delete[] arr;
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/array-shaping-5.C b/libgomp/testsuite/libgomp.c++/array-shaping-5.C
new file mode 100644
index 000000000000..7046a13c106f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/array-shaping-5.C
@@ -0,0 +1,38 @@
+// { dg-do run { target offload_device_nonshared_as } }
+
+#include <string.h>
+#include <assert.h>
+
+template<typename T, auto C>
+void foo (T *w, int e, int f, int g)
+{
+  memset (w, 0, sizeof (T) * 100);
+
+#pragma omp target enter data map(to: w[:100])
+
+  for (int j = 0; j < e; j++)
+    for (int i = 0; i < C; i++)
+      w[j * C + i] = i + j;
+
+#pragma omp target update to(([e][C]) w[3:2][f:g])
+
+#pragma omp target exit data map(from: w[:100])
+
+  for (int j = 0; j < e; j++)
+    for (int i = 0; i < C; i++)
+      if (j >= 3 && j < 5 && i >= f && i < f + g)
+	assert (w[j * C + i] == i + j);
+      else
+	assert (w[j * C + i] == 0.0f);
+}
+
+int main()
+{
+  float *arr = new float[100];
+
+  foo<float, 10> (arr, 10, 1, 8);
+
+  delete[] arr;
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/array-shaping-6.C b/libgomp/testsuite/libgomp.c++/array-shaping-6.C
new file mode 100644
index 000000000000..b960b5e58e14
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/array-shaping-6.C
@@ -0,0 +1,54 @@
+// { dg-do run { target offload_device_nonshared_as } }
+
+#include <assert.h>
+#include <string.h>
+
+template<typename T>
+void foo (T *&aref)
+{
+#pragma omp target enter data map(to: aref[:100])
+
+#pragma omp target
+  {
+    for (int i = 0; i < 10; i++)
+      for (int j = 0; j < 10; j++)
+	aref[i * 10 + j] = i + j;
+  }
+
+#pragma omp target update from(([10][10]) aref[2:3:2][7:3])
+
+  for (int i = 2; i < 8; i += 2)
+    for (int j = 7; j < 10; j++)
+      assert (aref[i * 10 + j] == i + j);
+
+#pragma omp target exit data map(delete: aref[:100])
+}
+
+int main()
+{
+  float *arr = new float[100];
+  float *&w = arr;
+
+  memset (arr, 0, 100 * sizeof (float));
+
+#pragma omp target enter data map(to: w[:100])
+
+#pragma omp target
+  {
+    for (int i = 0; i < 10; i++)
+      for (int j = 0; j < 10; j++)
+	w[i * 10 + j] = i + j;
+  }
+
+#pragma omp target update from(([10][10]) w[4:3][4:3])
+
+  for (int i = 4; i < 7; i++)
+    for (int j = 4; j < 7; j++)
+      assert (w[i * 10 + j] == i + j);
+
+#pragma omp target exit data map(delete: w[:100])
+
+  foo<float> (arr);
+
+  delete[] arr;
+}
diff --git a/libgomp/testsuite/libgomp.c++/array-shaping-7.C b/libgomp/testsuite/libgomp.c++/array-shaping-7.C
new file mode 100644
index 000000000000..b6193f8d619e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/array-shaping-7.C
@@ -0,0 +1,54 @@
+// { dg-do run { target offload_device_nonshared_as } }
+
+#include <assert.h>
+#include <string.h>
+
+template<typename T>
+void foo (T (&aref)[10][10])
+{
+#pragma omp target enter data map(to: aref)
+
+#pragma omp target
+  {
+    for (int i = 0; i < 10; i++)
+      for (int j = 0; j < 10; j++)
+	aref[i][j] = i + j;
+  }
+
+#pragma omp target update from(aref[2:3:2][7:3])
+
+  for (int i = 2; i < 8; i += 2)
+    for (int j = 7; j < 10; j++)
+      assert (aref[i][j] == i + j);
+
+#pragma omp target exit data map(delete: aref)
+}
+
+int main()
+{
+  float arr2d[10][10];
+  float (&w)[10][10] = arr2d;
+
+  memset (&arr2d, 0, 100 * sizeof (float));
+
+#pragma omp target enter data map(to: w)
+
+#pragma omp target
+  {
+    for (int i = 0; i < 10; i++)
+      for (int j = 0; j < 10; j++)
+	w[i][j] = i + j;
+  }
+
+#pragma omp target update from(w[4:3][4:3])
+
+  for (int i = 4; i < 7; i++)
+    for (int j = 4; j < 7; j++)
+      assert (w[i][j] == i + j);
+
+#pragma omp target exit data map(delete: w)
+
+  foo<float> (arr2d);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/array-shaping-8.C b/libgomp/testsuite/libgomp.c++/array-shaping-8.C
new file mode 100644
index 000000000000..a96cf3cffb80
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/array-shaping-8.C
@@ -0,0 +1,65 @@
+// { dg-do run { target offload_device_nonshared_as } }
+
+#include <assert.h>
+#include <string.h>
+
+template<typename T>
+struct C {
+  T *&aptr;
+
+  C(T *&aptr_1) : aptr(aptr_1)
+  {
+  }
+};
+
+template<typename T>
+void foo (T *c)
+{
+#pragma omp target enter data map(to: c->aptr, c->aptr[:100])
+
+#pragma omp target
+  {
+    for (int i = 0; i < 10; i++)
+      for (int j = 0; j < 10; j++)
+	c->aptr[i * 10 + j] = i + j;
+  }
+
+#pragma omp target update from(([10][10]) c->aptr[2:3:2][7:3])
+
+  for (int i = 2; i < 8; i += 2)
+    for (int j = 7; j < 10; j++)
+      assert (c->aptr[i * 10 + j] == i + j);
+
+#pragma omp target exit data map(delete: c->aptr, c->aptr[:100])
+}
+
+int main()
+{
+  float *arr = new float[100];
+  C<float> cvar(arr);
+
+  memset (arr, 0, 100 * sizeof (float));
+
+#pragma omp target enter data map(to: cvar.aptr, cvar.aptr[:100])
+
+#pragma omp target
+  {
+    for (int i = 0; i < 10; i++)
+      for (int j = 0; j < 10; j++)
+	cvar.aptr[i * 10 + j] = i + j;
+  }
+
+#pragma omp target update from(([10][10]) cvar.aptr[4:3][4:3])
+
+  for (int i = 4; i < 7; i++)
+    for (int j = 4; j < 7; j++)
+      assert (cvar.aptr[i * 10 + j] == i + j);
+
+#pragma omp target exit data map(delete: cvar.aptr, cvar.aptr[:100])
+
+  foo<C<float> > (&cvar);
+
+  delete[] arr;
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/array-shaping-9.C b/libgomp/testsuite/libgomp.c++/array-shaping-9.C
new file mode 100644
index 000000000000..786fe9d11edb
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/array-shaping-9.C
@@ -0,0 +1,95 @@
+// { dg-do run { target offload_device_nonshared_as } }
+
+#include <assert.h>
+#include <string.h>
+
+#define N 10
+
+struct B {
+  int (&aref)[N][N];
+
+  B(int (&aref1)[N][N]) : aref(aref1)
+  {
+  }
+};
+
+template<typename T, int S>
+struct C {
+  T (&aref)[S][S];
+
+  C(T (&aref1)[S][S]) : aref(aref1)
+  {
+  }
+};
+
+template<typename T>
+void foo (T *c)
+{
+#pragma omp target enter data map(to: c->aref)
+
+#pragma omp target
+  {
+    for (int i = 0; i < 10; i++)
+      for (int j = 0; j < 10; j++)
+	c->aref[i][j] = 2 * (i + j);
+  }
+
+#pragma omp target update from(c->aref[2:3:2][7:3])
+
+  for (int i = 2; i < 8; i += 2)
+    for (int j = 7; j < 10; j++)
+      assert (c->aref[i][j] == 2 * (i + j));
+
+#pragma omp target exit data map(delete: c->aref)
+}
+
+int main()
+{
+  int iarr[N][N];
+  float farr[N][N];
+  B bvar(iarr);
+  C<float, N> cvar(farr);
+
+  memset (iarr, 0, N * N * sizeof (int));
+  memset (farr, 0, N * N * sizeof (float));
+
+#pragma omp target enter data map(to: bvar.aref)
+
+#pragma omp target
+  {
+    for (int i = 0; i < 10; i++)
+      for (int j = 0; j < 10; j++)
+	bvar.aref[i][j] = i + j;
+  }
+
+#pragma omp target update from(bvar.aref[4:3][4:3])
+
+  for (int i = 4; i < 7; i++)
+    for (int j = 4; j < 7; j++)
+      assert (bvar.aref[i][j] == i + j);
+
+#pragma omp target exit data map(delete: bvar.aref)
+
+#pragma omp target enter data map(to: cvar.aref)
+
+#pragma omp target
+  {
+    for (int i = 0; i < 10; i++)
+      for (int j = 0; j < 10; j++)
+	cvar.aref[i][j] = i + j;
+  }
+
+#pragma omp target update from(cvar.aref[4:3][4:3])
+
+  for (int i = 4; i < 7; i++)
+    for (int j = 4; j < 7; j++)
+      assert (cvar.aref[i][j] == i + j);
+
+#pragma omp target exit data map(delete: cvar.aref)
+
+  memset (farr, 0, N * N * sizeof (float));
+
+  foo<C<float, N> > (&cvar);
+
+  return 0;
+}
