diff mbox series

[5/5] OpenMP: Array shaping operator and strided "target update" for C

Message ID 8e7811a8e3679d60fb804ac4df7e7523e0c27364.1688418868.git.julian@codesourcery.com
State New
Headers show
Series OpenMP: strides, rectangular updates and array-shaping operator for "target update" | expand

Commit Message

Julian Brown July 3, 2023, 9:33 p.m. UTC
Following the similar support for C++ and Fortran, here is the
C implementation for the OpenMP 5.0 array-shaping operator, and for
strided and rectangular updates for "target update".

Much of the implementation is shared with the C++ support added earlier
in this patch series.  Some details of parsing necessarily differ for C,
but the general ideas are the same.

2023-07-03  Julian Brown  <julian@codesourcery.com>

gcc/c/
	* c-parser.cc (c_parser_braced_init): Disallow array-shaping operator
	in braced init.
	(c_parser_conditional_expression): Disallow array-shaping operator in
	conditional expression.
	(c_parser_cast_expression): Add array-shaping operator support.
	(c_parser_postfix_expression): Disallow array-shaping operator in
	statement expressions.
	(c_parser_postfix_expression_after_primary): Add OpenMP array section
	stride support.
	(c_parser_expr_list): Disallow array-shaping operator in expression
	lists.
	(c_array_type_nelts_top, c_array_type_nelts_total): New functions.
	(c_parser_omp_variable_list): Support array-shaping operator.
	(c_parser_omp_target_update): Recognize GOMP_MAP_TO_GRID and
	GOMP_MAP_FROM_GRID map kinds as well as OMP_CLAUSE_TO/OMP_CLAUSE_FROM.
	* c-tree.h (c_omp_array_shaping_op_p, c_omp_has_array_shape_p): New
	extern declarations.
	(create_omp_arrayshape_type): Add prototype.
	* c-typeck.cc (c_omp_array_shaping_op_p, c_omp_has_array_shape_p): New
	globals.
	(build_omp_array_section): Permit integral types, not just integer
	constants, when creating array types for array sections.
	(create_omp_arrayshape_type): New function.
	(handle_omp_array_sections_1): Add DISCONTIGUOUS parameter.  Add
	strided/rectangular array section support.
	(omp_array_section_low_bound): New function.
	(handle_omp_array_sections): Add DISCONTIGUOUS parameter.  Add
	strided/rectangular array section support.
	(c_finish_omp_clauses): Update calls to handle_omp_array_sections.
	Handle discontiguous updates.

gcc/testsuite/
	* gcc.dg/gomp/bad-array-shaping-c-1.c: New test.
	* gcc.dg/gomp/bad-array-shaping-c-2.c: New test.
	* gcc.dg/gomp/bad-array-shaping-c-3.c: New test.
	* gcc.dg/gomp/bad-array-shaping-c-4.c: New test.
	* gcc.dg/gomp/bad-array-shaping-c-5.c: New test.
	* gcc.dg/gomp/bad-array-shaping-c-6.c: New test.
	* gcc.dg/gomp/bad-array-shaping-c-7.c: New test.

libgomp/
	* testsuite/libgomp.c/array-shaping-1.c: New test.
	* testsuite/libgomp.c/array-shaping-2.c: New test.
	* testsuite/libgomp.c/array-shaping-3.c: New test.
	* testsuite/libgomp.c/array-shaping-4.c: New test.
	* testsuite/libgomp.c/array-shaping-5.c: New test.
	* testsuite/libgomp.c/array-shaping-6.c: New test.
---
 gcc/c/c-parser.cc                             | 301 +++++++++++++++++-
 gcc/c/c-tree.h                                |   4 +
 gcc/c/c-typeck.cc                             | 241 ++++++++++++--
 .../gcc.dg/gomp/bad-array-shaping-c-1.c       |  26 ++
 .../gcc.dg/gomp/bad-array-shaping-c-2.c       |  24 ++
 .../gcc.dg/gomp/bad-array-shaping-c-3.c       |  30 ++
 .../gcc.dg/gomp/bad-array-shaping-c-4.c       |  27 ++
 .../gcc.dg/gomp/bad-array-shaping-c-5.c       |  17 +
 .../gcc.dg/gomp/bad-array-shaping-c-6.c       |  26 ++
 .../gcc.dg/gomp/bad-array-shaping-c-7.c       |  15 +
 libgomp/testsuite/libgomp.c/array-shaping-1.c | 236 ++++++++++++++
 libgomp/testsuite/libgomp.c/array-shaping-2.c |  39 +++
 libgomp/testsuite/libgomp.c/array-shaping-3.c |  42 +++
 libgomp/testsuite/libgomp.c/array-shaping-4.c |  36 +++
 libgomp/testsuite/libgomp.c/array-shaping-5.c |  38 +++
 libgomp/testsuite/libgomp.c/array-shaping-6.c |  45 +++
 16 files changed, 1099 insertions(+), 48 deletions(-)
 create mode 100644 gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-1.c
 create mode 100644 gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-2.c
 create mode 100644 gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-3.c
 create mode 100644 gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-4.c
 create mode 100644 gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-5.c
 create mode 100644 gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-6.c
 create mode 100644 gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-7.c
 create mode 100644 libgomp/testsuite/libgomp.c/array-shaping-1.c
 create mode 100644 libgomp/testsuite/libgomp.c/array-shaping-2.c
 create mode 100644 libgomp/testsuite/libgomp.c/array-shaping-3.c
 create mode 100644 libgomp/testsuite/libgomp.c/array-shaping-4.c
 create mode 100644 libgomp/testsuite/libgomp.c/array-shaping-5.c
 create mode 100644 libgomp/testsuite/libgomp.c/array-shaping-6.c
diff mbox series

Patch

diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 280426ddf10..7e895e11da2 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -5764,7 +5764,9 @@  c_parser_braced_init (c_parser *parser, tree type, bool nested_p,
   gcc_obstack_init (&braced_init_obstack);
   gcc_assert (c_parser_next_token_is (parser, CPP_OPEN_BRACE));
   bool save_c_omp_array_section_p = c_omp_array_section_p;
+  bool save_c_omp_array_shaping_op_p = c_omp_array_shaping_op_p;
   c_omp_array_section_p = false;
+  c_omp_array_shaping_op_p = false;
   matching_braces braces;
   braces.consume_open (parser);
   if (nested_p)
@@ -5804,6 +5806,7 @@  c_parser_braced_init (c_parser *parser, tree type, bool nested_p,
 	}
     }
   c_omp_array_section_p = save_c_omp_array_section_p;
+  c_omp_array_shaping_op_p = save_c_omp_array_shaping_op_p;
   c_token *next_tok = c_parser_peek_token (parser);
   if (next_tok->type != CPP_CLOSE_BRACE)
     {
@@ -8310,6 +8313,7 @@  c_parser_conditional_expression (c_parser *parser, struct c_expr *after,
   struct c_expr cond, exp1, exp2, ret;
   location_t start, cond_loc, colon_loc;
   bool save_c_omp_array_section_p = c_omp_array_section_p;
+  bool save_c_omp_array_shaping_op_p = c_omp_array_shaping_op_p;
 
   gcc_assert (!after || c_dialect_objc ());
 
@@ -8318,6 +8322,7 @@  c_parser_conditional_expression (c_parser *parser, struct c_expr *after,
   if (c_parser_next_token_is_not (parser, CPP_QUERY))
     return cond;
   c_omp_array_section_p = false;
+  c_omp_array_shaping_op_p = false;
   if (cond.value != error_mark_node)
     start = cond.get_start ();
   else
@@ -8371,6 +8376,7 @@  c_parser_conditional_expression (c_parser *parser, struct c_expr *after,
       ret.original_code = ERROR_MARK;
       ret.original_type = NULL;
       c_omp_array_section_p = save_c_omp_array_section_p;
+      c_omp_array_shaping_op_p = save_c_omp_array_shaping_op_p;
       return ret;
     }
   {
@@ -8418,6 +8424,7 @@  c_parser_conditional_expression (c_parser *parser, struct c_expr *after,
   set_c_expr_source_range (&ret, start, exp2.get_finish ());
   ret.m_decimal = 0;
   c_omp_array_section_p = save_c_omp_array_section_p;
+  c_omp_array_shaping_op_p = save_c_omp_array_shaping_op_p;
   return ret;
 }
 
@@ -8799,6 +8806,8 @@  c_parser_cast_expression (c_parser *parser, struct c_expr *after)
   if (after)
     return c_parser_postfix_expression_after_primary (parser,
 						      cast_loc, *after);
+  bool save_c_omp_has_array_shape_p = c_omp_has_array_shape_p;
+  c_omp_has_array_shape_p = false;
   /* If the expression begins with a parenthesized type name, it may
      be either a cast or a compound literal; we need to see whether
      the next character is '{' to tell the difference.  If not, it is
@@ -8807,6 +8816,10 @@  c_parser_cast_expression (c_parser *parser, struct c_expr *after)
   if (c_parser_next_token_is (parser, CPP_OPEN_PAREN)
       && c_token_starts_compound_literal (c_parser_peek_2nd_token (parser)))
     {
+      bool save_c_omp_array_section_p = c_omp_array_section_p;
+      bool save_c_omp_array_shaping_op_p = c_omp_array_shaping_op_p;
+      c_omp_array_section_p = false;
+      c_omp_array_shaping_op_p = false;
       struct c_declspecs *scspecs;
       struct c_type_name *type_name;
       struct c_expr ret;
@@ -8818,6 +8831,8 @@  c_parser_cast_expression (c_parser *parser, struct c_expr *after)
       parens.skip_until_found_close (parser);
       if (type_name == NULL)
 	{
+	  c_omp_array_section_p = save_c_omp_array_section_p;
+	  c_omp_array_shaping_op_p = save_c_omp_array_shaping_op_p;
 	  ret.set_error ();
 	  ret.original_code = ERROR_MARK;
 	  ret.original_type = NULL;
@@ -8828,9 +8843,15 @@  c_parser_cast_expression (c_parser *parser, struct c_expr *after)
       used_types_insert (type_name->specs->type);
 
       if (c_parser_next_token_is (parser, CPP_OPEN_BRACE))
-	return c_parser_postfix_expression_after_paren_type (parser, scspecs,
-							     type_name,
-							     cast_loc);
+	{
+	  c_expr r = c_parser_postfix_expression_after_paren_type (parser,
+								   scspecs,
+								   type_name,
+								   cast_loc);
+	  c_omp_array_section_p = save_c_omp_array_section_p;
+	  c_omp_array_shaping_op_p = save_c_omp_array_shaping_op_p;
+	  return r;
+	}
       if (scspecs)
 	error_at (cast_loc, "storage class specifier in cast");
       if (type_name->specs->alignas_p)
@@ -8847,10 +8868,61 @@  c_parser_cast_expression (c_parser *parser, struct c_expr *after)
       ret.original_code = ERROR_MARK;
       ret.original_type = NULL;
       ret.m_decimal = 0;
+      c_omp_array_section_p = save_c_omp_array_section_p;
+      c_omp_array_shaping_op_p = save_c_omp_array_shaping_op_p;
+      return ret;
+    }
+  else if (c_omp_array_shaping_op_p
+	   && c_parser_next_token_is (parser, CPP_OPEN_PAREN)
+	   && c_parser_peek_2nd_token (parser)->type == CPP_OPEN_SQUARE)
+    {
+      bool save_c_omp_array_section_p = c_omp_array_section_p;
+      bool save_c_omp_array_shaping_op_p = c_omp_array_shaping_op_p;
+      c_omp_array_section_p = false;
+      c_omp_array_shaping_op_p = false;
+      auto_vec<tree, 4> omp_shape_dims;
+      struct c_expr expr, ret;
+      matching_parens parens;
+      parens.consume_open (parser);
+      while (c_parser_next_token_is (parser, CPP_OPEN_SQUARE))
+	{
+	  c_parser_consume_token (parser);
+	  c_expr e = c_parser_expression (parser);
+	  if (e.value == error_mark_node)
+	    break;
+	  omp_shape_dims.safe_push (e.value);
+	  if (!c_parser_require (parser, CPP_CLOSE_SQUARE,
+				 "expected %<]%>"))
+	    break;
+	}
+      parens.require_close (parser);
+      c_omp_array_section_p = save_c_omp_array_section_p;
+      c_omp_array_shaping_op_p = save_c_omp_array_shaping_op_p;
+      {
+	location_t expr_loc = c_parser_peek_token (parser)->location;
+	bool save_c_omp_has_array_shape_p = c_omp_has_array_shape_p;
+	c_omp_has_array_shape_p = true;
+	expr = c_parser_cast_expression (parser, NULL);
+	c_omp_has_array_shape_p = save_c_omp_has_array_shape_p;
+	/* NOTE: We don't want to introduce conversions here.  */
+	expr = convert_lvalue_to_rvalue (expr_loc, expr, false, true);
+      }
+      tree arrtype
+	= create_omp_arrayshape_type (expr.value, &omp_shape_dims);
+      ret.value = build1_loc (cast_loc, VIEW_CONVERT_EXPR, arrtype,
+			      expr.value);
+      if (ret.value && expr.value)
+	set_c_expr_source_range (&ret, cast_loc, expr.get_finish ());
+      ret.original_code = ERROR_MARK;
+      ret.original_type = NULL;
+      ret.m_decimal = 0;
       return ret;
     }
   else
-    return c_parser_unary_expression (parser);
+    {
+      c_omp_has_array_shape_p = save_c_omp_has_array_shape_p;
+      return c_parser_unary_expression (parser);
+    }
 }
 
 /* Parse an unary expression (C90 6.3.3, C99 6.5.3, C11 6.5.3).
@@ -9860,6 +9932,7 @@  c_parser_postfix_expression (c_parser *parser)
 	  tree stmt;
 	  location_t brace_loc;
 	  bool save_c_omp_array_section_p = c_omp_array_section_p;
+	  bool save_c_omp_array_shaping_op_p = c_omp_array_shaping_op_p;
 	  c_parser_consume_token (parser);
 	  brace_loc = c_parser_peek_token (parser)->location;
 	  c_parser_consume_token (parser);
@@ -9877,6 +9950,7 @@  c_parser_postfix_expression (c_parser *parser)
 	      break;
 	    }
 	  c_omp_array_section_p = false;
+	  c_omp_array_shaping_op_p = false;
 	  stmt = c_begin_stmt_expr ();
 	  c_parser_compound_statement_nostart (parser);
 	  location_t close_loc = c_parser_peek_token (parser)->location;
@@ -9888,6 +9962,7 @@  c_parser_postfix_expression (c_parser *parser)
 	  set_c_expr_source_range (&expr, loc, close_loc);
 	  mark_exp_read (expr.value);
 	  c_omp_array_section_p = save_c_omp_array_section_p;
+	  c_omp_array_shaping_op_p = save_c_omp_array_shaping_op_p;
 	}
       else
 	{
@@ -11373,20 +11448,26 @@  c_parser_postfix_expression_after_primary (c_parser *parser,
 	  if (c_omp_array_section_p
 	      && c_parser_next_token_is (parser, CPP_COLON))
 	    {
+	      tree stride = NULL_TREE;
+
 	      c_parser_consume_token (parser);
 	      if (c_parser_next_token_is_not (parser, CPP_CLOSE_SQUARE))
 		len = c_parser_expression (parser).value;
 
+	      if (c_parser_next_token_is (parser, CPP_COLON))
+		{
+		  c_parser_consume_token (parser);
+		  if (c_parser_next_token_is_not (parser, CPP_CLOSE_SQUARE))
+		    stride = c_parser_expression (parser).value;
+		}
+
 	      c_parser_skip_until_found (parser, CPP_CLOSE_SQUARE,
 					 "expected %<]%>");
 
-	     /* NOTE: We are reusing using the type of the whole array as the
-		type of the array section here, which isn't necessarily
-		entirely correct.  Might need revisiting.  */
 	      start = expr.get_start ();
 	      finish = parser->tokens_buf[0].location;
 	      expr.value = build_omp_array_section (op_loc, expr.value, idx,
-						    len, NULL_TREE /* fixme */);
+						    len, stride);
 	      set_c_expr_source_range (&expr, start, finish);
 	      expr.original_code = ERROR_MARK;
 	      expr.original_type = NULL;
@@ -11397,7 +11478,20 @@  c_parser_postfix_expression_after_primary (c_parser *parser,
 					 "expected %<]%>");
 	      start = expr.get_start ();
 	      finish = parser->tokens_buf[0].location;
-	      expr.value = build_array_ref (op_loc, expr.value, idx);
+	      if (c_omp_has_array_shape_p)
+		/* 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
+		   c_parser_omp_variable_list).  */
+		expr.value
+		  = build_omp_array_section (op_loc, expr.value, idx,
+					     size_one_node, NULL_TREE);
+	      else
+		expr.value = build_array_ref (op_loc, expr.value, idx);
 	      set_c_expr_source_range (&expr, start, finish);
 	      expr.original_code = ERROR_MARK;
 	      expr.original_type = NULL;
@@ -11694,7 +11788,9 @@  c_parser_expr_list (c_parser *parser, bool convert_p, bool fold_p,
   struct c_expr expr;
   unsigned int idx = 0;
   bool save_c_omp_array_section_p = c_omp_array_section_p;
+  bool save_c_omp_array_shaping_op_p = c_omp_array_shaping_op_p;
   c_omp_array_section_p = false;
+  c_omp_array_shaping_op_p = false;
 
   ret = make_tree_vector ();
   if (p_orig_types == NULL)
@@ -11749,6 +11845,7 @@  c_parser_expr_list (c_parser *parser, bool convert_p, bool fold_p,
   if (orig_types)
     *p_orig_types = orig_types;
   c_omp_array_section_p = save_c_omp_array_section_p;
+  c_omp_array_shaping_op_p = save_c_omp_array_shaping_op_p;
   return ret;
 }
 
@@ -13939,6 +14036,35 @@  c_parser_oacc_wait_list (c_parser *parser, location_t clause_loc, tree list)
   return list;
 }
 
+/* Return, as an INTEGER_CST node, the number of elements for TYPE
+   (which is an ARRAY_TYPE).  This counts only elements of the top
+   array.  (From cp/tree.cc).  */
+
+static tree
+c_array_type_nelts_top (tree type)
+{
+  return fold_build2_loc (input_location, PLUS_EXPR, sizetype,
+			  array_type_nelts (type), size_one_node);
+}
+
+/* Return, as an INTEGER_CST node, the number of elements for TYPE
+   (which is an ARRAY_TYPE).  This one is a recursive count of all
+   ARRAY_TYPEs that are clumped together.  (From cp/tree.cc).  */
+
+static tree
+c_array_type_nelts_total (tree type)
+{
+  tree sz = c_array_type_nelts_top (type);
+  type = TREE_TYPE (type);
+  while (TREE_CODE (type) == ARRAY_TYPE)
+    {
+      tree n = c_array_type_nelts_top (type);
+      sz = fold_build2_loc (input_location, MULT_EXPR, sizetype, sz, n);
+      type = TREE_TYPE (type);
+    }
+  return sz;
+}
+
 /* OpenACC 2.0, OpenMP 2.5:
    variable-list:
      identifier
@@ -14067,12 +14193,24 @@  c_parser_omp_variable_list (c_parser *parser,
 	{
 	  location_t loc = c_parser_peek_token (parser)->location;
 	  bool save_c_omp_array_section_p = c_omp_array_section_p;
+	  bool save_c_omp_array_shaping_op_p = c_omp_array_shaping_op_p;
 	  c_omp_array_section_p = true;
+	  c_omp_array_shaping_op_p
+	    = (kind == OMP_CLAUSE_TO || kind == OMP_CLAUSE_FROM);
 	  c_expr expr = c_parser_expr_no_commas (parser, NULL);
 	  if (expr.value != error_mark_node)
 	    mark_exp_read (expr.value);
 	  c_omp_array_section_p = save_c_omp_array_section_p;
+	  c_omp_array_shaping_op_p = save_c_omp_array_shaping_op_p;
 	  tree decl = expr.value;
+	  tree reshaped_to = NULL_TREE;
+
+	  if (TREE_CODE (decl) == VIEW_CONVERT_EXPR
+	      && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
+	    {
+	      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
@@ -14085,6 +14223,31 @@  c_parser_omp_variable_list (c_parser *parser,
 	  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);
@@ -14093,18 +14256,63 @@  c_parser_omp_variable_list (c_parser *parser,
 		  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 (maybe_ptr_based_noncontig_update && sections == 0)
+		    break;
+
 		  if (TREE_CODE (decl) == COMPOUND_EXPR)
 		    {
 		      decl = TREE_OPERAND (decl, 1);
 		      STRIP_NOPS (decl);
 		    }
-		  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, NULL_TREE, loc,
@@ -14117,6 +14325,35 @@  c_parser_omp_variable_list (c_parser *parser,
 		      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.  */
+		      decl = build_indirect_ref (loc, decl, RO_UNARY_STAR);
+		      decl = build1_loc (loc, VIEW_CONVERT_EXPR, reshaped_to,
+					 decl);
 		    }
 		}
 
@@ -14144,6 +14381,14 @@  c_parser_omp_variable_list (c_parser *parser,
 	      decl = build_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 = c_array_type_nelts_total (reshaped_to);
+	      decl = build_omp_array_section (loc, decl, size_zero_node, elems,
+					      NULL_TREE);
+	    }
 	  else if (TREE_CODE (decl) == NON_LVALUE_EXPR
 		   || CONVERT_EXPR_P (decl))
 	    decl = TREE_OPERAND (decl, 0);
@@ -22777,8 +23022,38 @@  c_parser_omp_target_update (location_t loc, c_parser *parser,
   tree clauses
     = c_parser_omp_all_clauses (parser, OMP_TARGET_UPDATE_CLAUSE_MASK,
 				"#pragma omp target update");
-  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 (loc,
 		"%<#pragma omp target update%> must contain at least one "
diff --git a/gcc/c/c-tree.h b/gcc/c/c-tree.h
index 37790bab640..ee29f9de2cc 100644
--- a/gcc/c/c-tree.h
+++ b/gcc/c/c-tree.h
@@ -727,6 +727,8 @@  extern int in_sizeof;
 extern int in_typeof;
 extern bool c_in_omp_for;
 extern bool c_omp_array_section_p;
+extern bool c_omp_array_shaping_op_p;
+extern bool c_omp_has_array_shape_p;
 
 extern tree c_last_sizeof_arg;
 extern location_t c_last_sizeof_loc;
@@ -766,6 +768,8 @@  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, tree);
+extern tree create_omp_arrayshape_type (tree expr,
+					vec<tree> *omp_shape_dims);
 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 0eff41e7567..f6fb68cb491 100644
--- a/gcc/c/c-typeck.cc
+++ b/gcc/c/c-typeck.cc
@@ -79,6 +79,13 @@  bool c_in_omp_for;
 /* True when parsing OpenMP map clause.  */
 bool c_omp_array_section_p;
 
+/* True when parsing OpenMP to/from clause.  */
+bool c_omp_array_shaping_op_p;
+
+/* True if we have an OpenMP array-shaping "cast" expression.  This adjusts
+   the parsed representation for e.g. array refs.  */
+bool c_omp_has_array_shape_p;
+
 /* The argument of last parsed sizeof expression, only to be tested
    if expr.original_code == SIZEOF_EXPR.  */
 tree c_last_sizeof_arg;
@@ -2930,8 +2937,8 @@  build_omp_array_section (location_t loc, tree array, tree index, tree length,
 
   if (index != NULL_TREE
       && length != NULL_TREE
-      && TREE_CODE (index) == INTEGER_CST
-      && TREE_CODE (length) == INTEGER_CST)
+      && INTEGRAL_TYPE_P (TREE_TYPE (index))
+      && INTEGRAL_TYPE_P (TREE_TYPE (length)))
     {
       tree low = fold_convert (sizetype, index);
       tree high = fold_convert (sizetype, length);
@@ -2941,7 +2948,7 @@  build_omp_array_section (location_t loc, tree array, tree index, tree length,
     }
   else if ((index == NULL_TREE || integer_zerop (index))
 	   && length != NULL_TREE
-	   && TREE_CODE (length) == INTEGER_CST)
+	   && INTEGRAL_TYPE_P (TREE_TYPE (length)))
     idxtype = build_index_type (length);
   else
     idxtype = NULL_TREE;
@@ -2965,6 +2972,46 @@  build_omp_array_section (location_t loc, tree array, tree index, tree length,
 		     stride);
 }
 
+/* Build an array type whose dimensions are given by OMP_SHAPE_DIMS and whose
+   elements are of the type pointed to by the "base" node of EXPR with outer
+   OMP_ARRAY_SECTIONs and ARRAY_REFs stripped off, e.g. the type of "*myptr"
+   in "myptr[0:2:3][4][5:6]".  */
+
+tree
+create_omp_arrayshape_type (tree expr, vec<tree> *omp_shape_dims)
+{
+  tree strip_sections = expr;
+
+  while (TREE_CODE (strip_sections) == OMP_ARRAY_SECTION
+	 || TREE_CODE (strip_sections) == ARRAY_REF)
+    strip_sections = TREE_OPERAND (strip_sections, 0);
+
+  tree 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;
+  tree dim;
+  FOR_EACH_VEC_ELT_REVERSE (*omp_shape_dims, i, dim)
+    {
+      tree maxidx = fold_convert (sizetype, dim);
+      maxidx = size_binop (MINUS_EXPR, maxidx, size_one_node);
+      tree index = build_index_type (maxidx);
+      type = build_array_type (type, index);
+    }
+
+  return type;
+}
+
 
 /* Build an external reference to identifier ID.  FUN indicates
    whether this will be used for a function call.  LOC is the source
@@ -13714,7 +13761,8 @@  c_finish_omp_cancellation_point (location_t loc, tree clauses)
 static tree
 handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 			     bool &maybe_zero_len, unsigned int &first_non_one,
-			     bool &non_contiguous, enum c_omp_region_type ort)
+			     bool &non_contiguous, enum c_omp_region_type ort,
+			     int *discontiguous)
 {
   tree ret, low_bound, length, stride, type;
   bool openacc = (ort & C_ORT_ACC) != 0;
@@ -13795,11 +13843,14 @@  handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 
   ret = handle_omp_array_sections_1 (c, TREE_OPERAND (t, 0), types,
 				     maybe_zero_len, first_non_one,
-				     non_contiguous, ort);
+				     non_contiguous, 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);
@@ -13840,8 +13891,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))
@@ -13960,12 +14018,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 && !integer_onep (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)
@@ -13973,7 +14048,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))
 		    {
@@ -14047,7 +14122,8 @@  handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 	    {
 	      tree d_length = TREE_OPERAND (d, 2);
 	      tree d_stride = TREE_OPERAND (d, 3);
-	      if (d_length == NULL_TREE || !integer_onep (d_length)
+	      if (d_length == NULL_TREE
+		  || !integer_onep (d_length)
 		  || (d_stride && !integer_onep (d_stride)))
 		{
 		  if (openacc && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
@@ -14068,10 +14144,15 @@  handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 		      return error_mark_node;
 		    }
 
-		  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;
+		    }
 		}
 	    }
 	}
@@ -14083,7 +14164,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 = save_expr (low_bound);
   if (lb != low_bound)
@@ -14091,14 +14172,42 @@  handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
       TREE_OPERAND (t, 1) = lb;
       low_bound = lb;
     }
-  ret = build_array_ref (OMP_CLAUSE_LOCATION (c), ret, low_bound);
+  /* 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 = build_omp_array_section (OMP_CLAUSE_LOCATION (c), ret, low_bound,
+				   length, stride);
+  else
+    ret = build_array_ref (OMP_CLAUSE_LOCATION (c), ret, low_bound);
   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 build_array_ref (loc, ret, low_bound);
+    }
+
+  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.  */
 
 static bool
-handle_omp_array_sections (tree *pc, enum c_omp_region_type ort)
+handle_omp_array_sections (tree *pc, enum c_omp_region_type ort,
+			   int *discontiguous)
 {
   tree c = *pc;
   bool maybe_zero_len = false;
@@ -14114,7 +14223,7 @@  handle_omp_array_sections (tree *pc, 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,
-					    non_contiguous, ort);
+					    non_contiguous, ort, discontiguous);
   if (first == error_mark_node)
     return true;
   if (first == NULL_TREE)
@@ -14153,11 +14262,14 @@  handle_omp_array_sections (tree *pc, enum c_omp_region_type ort)
       if (int_size_in_bytes (TREE_TYPE (first)) <= 0)
 	maybe_zero_len = true;
 
+      bool higher_discontiguous = false;
+
       for (i = num, t = OMP_CLAUSE_DECL (c); i > 0;
 	   t = TREE_OPERAND (t, 0))
 	{
 	  tree low_bound = TREE_OPERAND (t, 1);
 	  tree length = TREE_OPERAND (t, 2);
+	  tree stride = TREE_OPERAND (t, 3);
 
 	  i--;
 	  if (low_bound
@@ -14170,6 +14282,11 @@  handle_omp_array_sections (tree *pc, 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;
 
@@ -14179,10 +14296,49 @@  handle_omp_array_sections (tree *pc, enum c_omp_region_type ort)
 	      continue;
 	    }
 
+	  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])
@@ -14196,12 +14352,17 @@  handle_omp_array_sections (tree *pc, 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 (length != NULL_TREE
@@ -14321,6 +14482,8 @@  handle_omp_array_sections (tree *pc, 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);
       first = c_fully_fold (first, false, NULL);
       OMP_CLAUSE_DECL (c) = first;
       if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
@@ -14329,7 +14492,8 @@  handle_omp_array_sections (tree *pc, enum c_omp_region_type ort)
 	size = c_fully_fold (size, false, NULL);
       OMP_CLAUSE_SIZE (c) = size;
 
-      if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+      if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+	  && !(discontiguous && *discontiguous == 2))
 	return false;
 
       auto_vec<omp_addr_token *, 10> addr_tokens;
@@ -14342,7 +14506,8 @@  handle_omp_array_sections (tree *pc, enum c_omp_region_type ort)
       tree *npc = ai.expand_map_clause (pc, first, addr_tokens, ort);
       if (npc != NULL)
 	{
-	  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;
 
 	  return false;
@@ -14689,7 +14854,7 @@  c_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, ort))
+	      if (handle_omp_array_sections (pc, ort, NULL))
 		{
 		  remove = true;
 		  break;
@@ -15428,7 +15593,7 @@  c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	    last_iterators = NULL_TREE;
 	  if (TREE_CODE (t) == OMP_ARRAY_SECTION)
 	    {
-	      if (handle_omp_array_sections (pc, ort))
+	      if (handle_omp_array_sections (pc, ort, NULL))
 		remove = true;
 	      else if ((c = *pc)
 		       && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND
@@ -15535,6 +15700,9 @@  c_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:
@@ -15549,7 +15717,10 @@  c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		grp_start_p = pc;
 		grp_sentinel = OMP_CLAUSE_CHAIN (c);
 
-		if (handle_omp_array_sections (pc, ort))
+		int discontiguous
+		  = (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO
+		     || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM);
+		if (handle_omp_array_sections (pc, ort, &discontiguous))
 		  remove = true;
 		else
 		  {
@@ -15946,7 +16117,7 @@  c_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, ort))
+	      if (handle_omp_array_sections (pc, ort, NULL))
 		remove = true;
 	      else
 		{
diff --git a/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-1.c b/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-1.c
new file mode 100644
index 00000000000..42d584fa624
--- /dev/null
+++ b/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-1.c
@@ -0,0 +1,26 @@ 
+// { dg-do compile }
+
+#include <string.h>
+#include <assert.h>
+#include <stdlib.h>
+
+int main (void)
+{
+  float *arr = calloc (100, sizeof (float));
+
+#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])
+
+  free (arr);
+
+  return 0;
+}
diff --git a/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-2.c b/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-2.c
new file mode 100644
index 00000000000..6be3e009ecb
--- /dev/null
+++ b/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-2.c
@@ -0,0 +1,24 @@ 
+// { dg-do compile }
+
+#include <string.h>
+#include <assert.h>
+#include <stdlib.h>
+
+int main (void)
+{
+  float *arr = calloc (100, sizeof (float));
+
+  /* This isn't allowed.  */
+#pragma omp target enter data map(to: ([10][10]) arr[:100])
+/* { dg-error {expected expression before '\[' token} "" { target *-*-* } .-1 } */
+/* { dg-error {'#pragma omp target enter data' must contain at least one 'map' clause} "" { target *-*-* } .-2 } */
+
+  /* Nor this.  */
+#pragma omp target exit data map(from: ([10][10]) arr[:100])
+/* { dg-error {expected expression before '\[' token} "" { target *-*-* } .-1 } */
+/* { dg-error {'#pragma omp target exit data' must contain at least one 'map' clause} "" { target *-*-* } .-2 } */
+
+  free (arr);
+
+  return 0;
+}
diff --git a/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-3.c b/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-3.c
new file mode 100644
index 00000000000..1715b8ff9ed
--- /dev/null
+++ b/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-3.c
@@ -0,0 +1,30 @@ 
+// { dg-do compile }
+
+#include <string.h>
+#include <assert.h>
+#include <stdlib.h>
+
+extern float* baz(void*);
+
+int main (void)
+{
+  float *arr = calloc (100, sizeof (float));
+  int c = 50;
+
+#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;
+
+  /* No array shaping inside a function call.  */
+#pragma omp target update to(baz(([10][10]) arr))
+/* { dg-error {expected expression before '\[' token} "" { 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])
+
+  free (arr);
+
+  return 0;
+}
diff --git a/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-4.c b/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-4.c
new file mode 100644
index 00000000000..cebefd36d18
--- /dev/null
+++ b/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-4.c
@@ -0,0 +1,27 @@ 
+// { dg-do compile }
+
+#include <string.h>
+#include <assert.h>
+#include <stdlib.h>
+
+int main (void)
+{
+  float *arr = calloc (100, sizeof (float));
+
+#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;
+
+  /* No array shaping inside a statement expression.  */
+#pragma omp target update to( ({ int d = 10; ([d][d]) arr; }) )
+/* { dg-error {expected expression before '\[' token} "" { 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])
+
+  free (arr);
+
+  return 0;
+}
diff --git a/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-5.c b/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-5.c
new file mode 100644
index 00000000000..e1c4991f5c3
--- /dev/null
+++ b/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-5.c
@@ -0,0 +1,17 @@ 
+// { dg-do compile }
+
+struct S {
+  void *pp;
+};
+
+int main()
+{
+  int *sub1;
+
+  /* No array section inside compound literal.  */
+#pragma omp target update to( (struct S) { .pp = ([10][10]) sub1 } )
+/* { dg-error {expected expression before '\[' token} "" { target *-*-* } .-1 } */
+/* { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-2 } */
+
+  return 0;
+}
diff --git a/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-6.c b/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-6.c
new file mode 100644
index 00000000000..d282d8598b2
--- /dev/null
+++ b/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-6.c
@@ -0,0 +1,26 @@ 
+// { dg-do compile }
+
+int main (void)
+{
+  char *ptr;
+
+#pragma omp target update to(([5][6][7]) ptr[0:4][0:7][0:7])
+/* { dg-error {length '7' 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 } */
+
+  return 0;
+}
diff --git a/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-7.c b/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-7.c
new file mode 100644
index 00000000000..233d8da6f44
--- /dev/null
+++ b/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-7.c
@@ -0,0 +1,15 @@ 
+/* { dg-do compile } */
+
+int cond;
+
+int main (void)
+{
+  int *arr;
+
+  /* No array shaping inside conditional operator.  */
+#pragma omp target update to(cond ? ([3][9]) arr : ([2][7]) arr)
+/* { dg-error {expected expression before '\[' token} "" { target *-*-* } .-1 } */
+/* { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-2 } */
+
+  return 0;
+}
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 00000000000..808c5f9ceae
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/array-shaping-1.c
@@ -0,0 +1,236 @@ 
+// { dg-do run { target offload_device_nonshared_as } }
+
+#include <string.h>
+#include <assert.h>
+#include <stdlib.h>
+
+volatile int yy = 4, zz = 2, str_str = 2;
+
+int main()
+{
+  int *arr;
+  int x = 5;
+  int arr2d[10][10];
+
+  arr = calloc (100, sizeof (int));
+
+  /* Update whole reshaped array.  */
+
+#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);
+
+  free (arr);
+
+  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 00000000000..42a6e0ca7d8
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/array-shaping-2.c
@@ -0,0 +1,39 @@ 
+// { dg-do run { target offload_device_nonshared_as } }
+
+#include <assert.h>
+#include <stdlib.h>
+
+typedef struct {
+  int *aptr;
+} C;
+
+int main()
+{
+  C cvar;
+
+  cvar.aptr = calloc (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 = 0; i < 10; i++)
+    for (int j = 0; j < 10; j++)
+      if (i >= 4 && i < 7 && j >= 4 && j < 7)
+	assert (cvar.aptr[i * 10 + j] == i + j);
+      else
+	assert (cvar.aptr[i * 10 + j] == 0);
+
+#pragma omp target exit data map(delete: cvar.aptr, cvar.aptr[:100])
+
+  free (cvar.aptr);
+
+  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 00000000000..5dda2e32832
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/array-shaping-3.c
@@ -0,0 +1,42 @@ 
+// { dg-do run { target offload_device_nonshared_as } }
+
+#include <assert.h>
+#include <stdlib.h>
+#include <string.h>
+
+#define N 10
+
+typedef struct {
+  int arr[N][N];
+} B;
+
+int main()
+{
+  B *bvar = malloc (sizeof (B));
+
+  memset (bvar, 0, sizeof (B));
+
+#pragma omp target enter data map(to: bvar->arr)
+
+#pragma omp target
+  {
+    for (int i = 0; i < 10; i++)
+      for (int j = 0; j < 10; j++)
+	bvar->arr[i][j] = i + j;
+  }
+
+#pragma omp target update from(bvar->arr[4:3][4:3])
+
+  for (int i = 0; i < 10; i++)
+    for (int j = 0; j < 10; j++)
+      if (i >= 4 && i < 7 && j >= 4 && j < 7)
+	assert (bvar->arr[i][j] == i + j);
+      else
+	assert (bvar->arr[i][j] == 0);
+
+#pragma omp target exit data map(delete: bvar->arr)
+
+  free (bvar);
+
+  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 00000000000..2b9e6949b60
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/array-shaping-4.c
@@ -0,0 +1,36 @@ 
+// { dg-do run { target offload_device_nonshared_as } }
+
+#include <assert.h>
+#include <string.h>
+
+#define N 10
+
+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 = 0; i < 10; i++)
+    for (int j = 0; j < 10; j++)
+      if (i >= 4 && i < 7 && j >= 4 && j < 7)
+	assert (iarr[i * 10 + j] == i + j);
+      else
+	assert (iarr[i * 10 + j] == 0);
+
+#pragma omp target exit data map(delete: iarr)
+
+  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 00000000000..1034682e4ca
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/array-shaping-5.c
@@ -0,0 +1,38 @@ 
+// { dg-do run { target offload_device_nonshared_as } }
+
+#include <assert.h>
+#include <string.h>
+
+#define N 10
+
+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 = 0; i < 10; i++)
+    for (int j = 0; j < 10; j++)
+      if (i >= 4 && i < 7 && j >= 4 && j < 7)
+	assert (iarrp[i * 10 + j] == i + j);
+      else
+	assert (iarrp[i * 10 + j] == 0);
+
+#pragma omp target exit data map(delete: iarr_real)
+
+  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 00000000000..59388232244
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/array-shaping-6.c
@@ -0,0 +1,45 @@ 
+// { dg-do run { target offload_device_nonshared_as } }
+
+#include <assert.h>
+#include <stdlib.h>
+#include <string.h>
+
+#define N 10
+
+int main ()
+{
+  int *iptr = calloc (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);
+
+  free (iptr);
+
+  return 0;
+}