Non-contiguous array support patches [PR76739]

This is based on OG14 commit b143c1c447945ce05903ff1360ead97774dfce4b,
which was based from v4, posted upstream here:
https://gcc.gnu.org/pipermail/gcc-patches/2020-April/543437.html

It also incorporates a number of follow-up bug and bit-rot fixes, OG14
commits
e11726d3467543de45448097dde27ba34bf04bfe
87ea4de1c4a360d5d62357491a41811213f4528c
151fc161d0ed640048444ca18f9325e3d2e03e99
628a000bdbf63252c2ede13ccab8e99a19769866
11263c048d39ab1d6a11067b18674bf8307bbbf5
8c1068bbe3e52529bede5466a43af8d98f38dac2

gcc/c/ChangeLog
	PR other/76739
	* c-typeck.cc (handle_omp_array_sections_1): Add 'bool &non_contiguous'
	parameter, adjust recursive call site, add cases for allowing
	pointer based multi-dimensional arrays for OpenACC.  Reject
	non-DECL base-pointer cases as unsupported.
	(handle_omp_array_sections): Adjust handle_omp_array_sections_1 call,
	handle non-contiguous case to create dynamic array map.

gcc/cp/ChangeLog
	PR other/76739
	* semantics.cc (handle_omp_array_sections_1): Add 'bool &non_contiguous'
	parameter, adjust recursive call site, add cases for allowing
	pointer based multi-dimensional arrays for OpenACC.  Reject
	non-DECL base-pointer cases as unsupported.
	(handle_omp_array_sections): Adjust handle_omp_array_sections_1 call,
	handle non-contiguous case to create dynamic array map.

gcc/fortran/ChangeLog
	PR other/76739
	* f95-lang.cc (DEF_FUNCTION_TYPE_VAR_5): New symbol.
	* types.def (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_VAR): New type.

gcc/ChangeLog
	PR other/76739
	* builtin-types.def (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_VAR): New type.
	* gimplify.cc (omp_group_base): Handle GOMP_MAP_NONCONTIG_ARRAY_*.
	(gimplify_scan_omp_clauses): Handle OMP_TARGET_UPDATE.
	(gimplify_adjust_omp_clauses): Skip gimplification of
	OMP_CLAUSE_SIZE of non-contiguous array maps (which is a TREE_LIST).
	* omp-builtins.def (BUILT_IN_GOACC_DATA_START): Adjust function type
	to new BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_VAR.
	* omp-expand.cc (expand_omp_target): Add non-contiguous array
	descriptor pointers to variadic arguments.
	* omp-low.cc (append_field_to_record_type): New function.
	(create_noncontig_array_descr_type): Likewise.
	(create_noncontig_array_descr_init_code): Likewise.
	(scan_sharing_clauses): For non-contiguous array map kinds, check for
	supported dimension structure, and install non-contiguous array
	variable into current omp_context.
	(reorder_noncontig_array_clauses): New function.
	(scan_omp_target): Call reorder_noncontig_array_clauses to place
	non-contiguous array map clauses at beginning of clause sequence.
	(lower_omp_target): Add handling for non-contiguous array map kinds,
	add all created non-contiguous array descriptors to
	gimple_omp_target_data_arg.
	* tree-pretty-print.cc (dump_omp_clause): Handle
	GOMP_MAP_NONCONTIG_ARRAY_*.

gcc/testsuite/ChangeLog
	PR other/76739
	* c-c++-common/goacc/data-clause-1.c (foo): Remove expected message.
	* c-c++-common/goacc/noncontig_array-1.c: New test.
	* g++.dg/goacc/data-clause-1.C (foo): Remove expected message.

include/ChangeLog
	PR other/76739
	* gomp-constants.h (GOMP_MAP_FLAG_SPECIAL_3): Define.
	(enum gomp_map_kind): Add GOMP_MAP_NONCONTIG_ARRAY,
	GOMP_MAP_NONCONTIG_ARRAY_TO, GOMP_MAP_NONCONTIG_ARRAY_FROM,
	GOMP_MAP_NONCONTIG_ARRAY_TOFROM, GOMP_MAP_NONCONTIG_ARRAY_FORCE_TO,
	GOMP_MAP_NONCONTIG_ARRAY_FORCE_FROM,
	GOMP_MAP_NONCONTIG_ARRAY_FORCE_TOFROM,
	GOMP_MAP_NONCONTIG_ARRAY_ALLOC, GOMP_MAP_NONCONTIG_ARRAY_FORCE_ALLOC,
	GOMP_MAP_NONCONTIG_ARRAY_FORCE_PRESENT.
	(GOMP_MAP_NONCONTIG_ARRAY_P): Define.

libgomp/ChangeLog
	PR other/76739
	* libgomp.h (gomp_map_vars_openacc): New function declaration.
	* libgomp_g.h (GOACC_data_start): Add variadic '...' to declaration.
	* oacc-int.h (struct goacc_ncarray_dim): New struct declaration.
	(struct goacc_ncarray_descr_type): Likewise.
	(struct goacc_ncarray): Likewise.
	(struct goacc_ncarray_info): Likewise.
	(goacc_noncontig_array_create_ptrblock): New function declaration.
	* oacc-parallel.c (goacc_noncontig_array_count_rows): New function.
	(goacc_noncontig_array_compute_sizes): Likewise.
	(goacc_noncontig_array_fill_rows_1): Likewise.
	(goacc_noncontig_array_fill_rows): Likewise.
	(goacc_process_noncontiguous_arrays): Likewise.
	(goacc_noncontig_array_create_ptrblock): Likewise.
	(GOACC_parallel_keyed): Use goacc_process_noncontiguous_arrays to
	handle non-contiguous array descriptors at end of varargs, adjust
	to use gomp_map_vars_openacc.
	(GOACC_data_start): Likewise. Adjust function type to accept varargs.
	* target.c (gomp_map_vars_internal): Add struct goacc_ncarray_info *
	nca_info parameter, add handling code for non-contiguous arrays.
	(gomp_map_vars_openacc): Add new function for specialization of
	gomp_map_vars_internal for OpenACC structured region usage.
	* testsuite/libgomp.oacc-c-c++-common/noncontig_array-1.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/noncontig_array-2.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/noncontig_array-3.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/noncontig_array-4.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/noncontig_array-utils.h: Support
	header for new tests.

Co-Authored-By: Kwok Cheung Yeung <kcy@codesourcery.com>
Co-Authored-By: Paul-Antoine Arras <parras@baylibre.com>

13 files changed, 460 insertions, 16 deletions

diff --git a/gcc/builtin-types.def b/gcc/builtin-types.def
index 9583d30..88bf917 100644
--- a/gcc/builtin-types.def
+++ b/gcc/builtin-types.def
@@ -1073,6 +1073,9 @@ DEF_FUNCTION_TYPE_VAR_5 (BT_FN_INT_STRING_SIZE_INT_SIZE_CONST_STRING_VAR,
 DEF_FUNCTION_TYPE_VAR_5 (BT_FN_INT_INT_INT_INT_INT_INT_VAR,
 			 BT_INT, BT_INT, BT_INT, BT_INT, BT_INT, BT_INT)
 
+DEF_FUNCTION_TYPE_VAR_5 (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_VAR,
+			 BT_VOID, BT_INT, BT_SIZE, BT_PTR, BT_PTR, BT_PTR)
+
 DEF_FUNCTION_TYPE_VAR_6 (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_VAR,
 			 BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_SIZE,
 			 BT_PTR, BT_PTR, BT_PTR)
diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc
index ea83451..a7ff2a5 100644
--- a/gcc/c/c-typeck.cc
+++ b/gcc/c/c-typeck.cc
@@ -15172,12 +15172,14 @@ c_finish_omp_cancellation_point (location_t loc, tree clauses)
    <= FIRST_NON_ONE we diagnose non-contiguous arrays if low bound isn't
    0 or length isn't the array domain max + 1, for > FIRST_NON_ONE we
    can if MAYBE_ZERO_LEN is false.  MAYBE_ZERO_LEN will be true in the above
-   case though, as some lengths could be zero.  */
+   case though, as some lengths could be zero.
+   NON_CONTIGUOUS will be true if this is an OpenACC non-contiguous array
+   section.  */
 
 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)
+			     bool &non_contiguous, enum c_omp_region_type ort)
 {
   tree ret, low_bound, length, type;
   bool openacc = (ort & C_ORT_ACC) != 0;
@@ -15257,7 +15259,8 @@ 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, ort);
+				     maybe_zero_len, first_non_one,
+				     non_contiguous, ort);
   if (ret == error_mark_node || ret == NULL_TREE)
     return ret;
 
@@ -15484,7 +15487,9 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 	  return error_mark_node;
 	}
       /* If there is a pointer type anywhere but in the very first
-	 array-section-subscript, the array section could be non-contiguous.  */
+	 array-section-subscript, the array section could be non-contiguous.
+	 Note that OpenACC does accept these kinds of non-contiguous pointer
+	 based arrays.  */
       if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND
 	  && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_AFFINITY
 	  && TREE_CODE (TREE_OPERAND (t, 0)) == OMP_ARRAY_SECTION)
@@ -15498,6 +15503,24 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 	      tree d_length = TREE_OPERAND (d, 2);
 	      if (d_length == NULL_TREE || !integer_onep (d_length))
 		{
+		  if (openacc && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
+		    {
+		      while (TREE_CODE (d) == OMP_ARRAY_SECTION)
+			d = TREE_OPERAND (d, 0);
+		      if (DECL_P (d))
+			{
+			  /* Note that OpenACC does accept these kinds of
+			     non-contiguous pointer based arrays.  */
+			  non_contiguous = true;
+			  break;
+			}
+		      error_at (OMP_CLAUSE_LOCATION (c),
+				"base-pointer expression in %qs clause not "
+				"supported for non-contiguous arrays",
+				omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+		      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)]);
@@ -15532,6 +15555,7 @@ handle_omp_array_sections (tree &c, enum c_omp_region_type ort)
 {
   bool maybe_zero_len = false;
   unsigned int first_non_one = 0;
+  bool non_contiguous = false;
   auto_vec<tree, 10> types;
   tree *tp = &OMP_CLAUSE_DECL (c);
   if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND
@@ -15542,7 +15566,7 @@ handle_omp_array_sections (tree &c, 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);
+					    non_contiguous, ort);
   if (first == error_mark_node)
     return true;
   if (first == NULL_TREE)
@@ -15576,6 +15600,7 @@ handle_omp_array_sections (tree &c, enum c_omp_region_type ort)
       unsigned int num = types.length (), i;
       tree t, side_effects = NULL_TREE, size = NULL_TREE;
       tree condition = NULL_TREE;
+      tree ncarray_dims = NULL_TREE;
 
       if (int_size_in_bytes (TREE_TYPE (first)) <= 0)
 	maybe_zero_len = true;
@@ -15599,6 +15624,13 @@ handle_omp_array_sections (tree &c, enum c_omp_region_type ort)
 	    length = fold_convert (sizetype, length);
 	  if (low_bound == NULL_TREE)
 	    low_bound = integer_zero_node;
+
+	  if (non_contiguous)
+	    {
+	      ncarray_dims = tree_cons (low_bound, length, ncarray_dims);
+	      continue;
+	    }
+
 	  if (!maybe_zero_len && i > first_non_one)
 	    {
 	      if (integer_nonzerop (low_bound))
@@ -15695,6 +15727,14 @@ handle_omp_array_sections (tree &c, enum c_omp_region_type ort)
 		size = size_binop (MULT_EXPR, size, l);
 	    }
 	}
+      if (non_contiguous)
+	{
+	  int kind = OMP_CLAUSE_MAP_KIND (c);
+	  OMP_CLAUSE_SET_MAP_KIND (c, kind | GOMP_MAP_NONCONTIG_ARRAY);
+	  OMP_CLAUSE_DECL (c) = t;
+	  OMP_CLAUSE_SIZE (c) = ncarray_dims;
+	  return false;
+	}
       if (side_effects)
 	size = build2 (COMPOUND_EXPR, sizetype, side_effects, size);
       if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION
diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc
index a10ef34..d0bde20 100644
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -5906,12 +5906,14 @@ public:
    <= FIRST_NON_ONE we diagnose non-contiguous arrays if low bound isn't
    0 or length isn't the array domain max + 1, for > FIRST_NON_ONE we
    can if MAYBE_ZERO_LEN is false.  MAYBE_ZERO_LEN will be true in the above
-   case though, as some lengths could be zero.  */
+   case though, as some lengths could be zero.
+   NON_CONTIGUOUS will be true if this is an OpenACC non-contiguous array
+   section.  */
 
 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)
+			     bool &non_contiguous, enum c_omp_region_type ort)
 {
   tree ret, low_bound, length, type;
   bool openacc = (ort & C_ORT_ACC) != 0;
@@ -5975,7 +5977,8 @@ 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,
+				     non_contiguous, ort);
   if (ret == error_mark_node || ret == NULL_TREE)
     return ret;
 
@@ -6228,6 +6231,24 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 	      tree d_length = TREE_OPERAND (d, 2);
 	      if (d_length == NULL_TREE || !integer_onep (d_length))
 		{
+		  if (openacc && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
+		    {
+		      while (TREE_CODE (d) == OMP_ARRAY_SECTION)
+			d = TREE_OPERAND (d, 0);
+		      if (DECL_P (d))
+			{
+			  /* Note that OpenACC does accept these kinds of
+			     non-contiguous pointer based arrays.  */
+			  non_contiguous = true;
+			  break;
+			}
+		      error_at (OMP_CLAUSE_LOCATION (c),
+				"base-pointer expression in %qs clause not "
+				"supported for non-contiguous arrays",
+				omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+		      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)]);
@@ -6274,6 +6295,7 @@ handle_omp_array_sections (tree &c, enum c_omp_region_type ort)
 {
   bool maybe_zero_len = false;
   unsigned int first_non_one = 0;
+  bool non_contiguous = false;
   auto_vec<tree, 10> types;
   tree *tp = &OMP_CLAUSE_DECL (c);
   if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND
@@ -6284,7 +6306,7 @@ handle_omp_array_sections (tree &c, 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);
+					    non_contiguous, ort);
   if (first == error_mark_node)
     return true;
   if (first == NULL_TREE)
@@ -6319,6 +6341,7 @@ handle_omp_array_sections (tree &c, enum c_omp_region_type ort)
       unsigned int num = types.length (), i;
       tree t, side_effects = NULL_TREE, size = NULL_TREE;
       tree condition = NULL_TREE;
+      tree ncarray_dims = NULL_TREE;
 
       if (int_size_in_bytes (TREE_TYPE (first)) <= 0)
 	maybe_zero_len = true;
@@ -6346,6 +6369,13 @@ handle_omp_array_sections (tree &c, enum c_omp_region_type ort)
 	    length = fold_convert (sizetype, length);
 	  if (low_bound == NULL_TREE)
 	    low_bound = integer_zero_node;
+
+	  if (non_contiguous)
+	    {
+	      ncarray_dims = tree_cons (low_bound, length, ncarray_dims);
+	      continue;
+	    }
+
 	  if (!maybe_zero_len && i > first_non_one)
 	    {
 	      if (integer_nonzerop (low_bound))
@@ -6437,6 +6467,14 @@ handle_omp_array_sections (tree &c, enum c_omp_region_type ort)
 	}
       if (!processing_template_decl)
 	{
+	  if (non_contiguous)
+	    {
+	      int kind = OMP_CLAUSE_MAP_KIND (c);
+	      OMP_CLAUSE_SET_MAP_KIND (c, kind | GOMP_MAP_NONCONTIG_ARRAY);
+	      OMP_CLAUSE_DECL (c) = t;
+	      OMP_CLAUSE_SIZE (c) = ncarray_dims;
+	      return false;
+	    }
 	  if (side_effects)
 	    size = build2 (COMPOUND_EXPR, sizetype, side_effects, size);
 	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION
diff --git a/gcc/fortran/f95-lang.cc b/gcc/fortran/f95-lang.cc
index 1f09553..21f0690 100644
--- a/gcc/fortran/f95-lang.cc
+++ b/gcc/fortran/f95-lang.cc
@@ -674,6 +674,8 @@ gfc_init_builtin_functions (void)
 #define DEF_FUNCTION_TYPE_VAR_0(NAME, RETURN) NAME,
 #define DEF_FUNCTION_TYPE_VAR_1(NAME, RETURN, ARG1) NAME,
 #define DEF_FUNCTION_TYPE_VAR_2(NAME, RETURN, ARG1, ARG2) NAME,
+#define DEF_FUNCTION_TYPE_VAR_5(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5) \
+				 NAME,
 #define DEF_FUNCTION_TYPE_VAR_6(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
 				 ARG6) NAME,
 #define DEF_FUNCTION_TYPE_VAR_7(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
@@ -696,6 +698,7 @@ gfc_init_builtin_functions (void)
 #undef DEF_FUNCTION_TYPE_VAR_0
 #undef DEF_FUNCTION_TYPE_VAR_1
 #undef DEF_FUNCTION_TYPE_VAR_2
+#undef DEF_FUNCTION_TYPE_VAR_5
 #undef DEF_FUNCTION_TYPE_VAR_6
 #undef DEF_FUNCTION_TYPE_VAR_7
 #undef DEF_POINTER_TYPE
@@ -1208,6 +1211,15 @@ gfc_init_builtin_functions (void)
 					builtin_types[(int) ARG1],     	\
 					builtin_types[(int) ARG2],     	\
 					NULL_TREE);
+#define DEF_FUNCTION_TYPE_VAR_5(ENUM, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5) \
+  builtin_types[(int) ENUM]						\
+    = build_varargs_function_type_list (builtin_types[(int) RETURN],   	\
+					builtin_types[(int) ARG1],     	\
+					builtin_types[(int) ARG2],     	\
+					builtin_types[(int) ARG3],	\
+					builtin_types[(int) ARG4],	\
+					builtin_types[(int) ARG5],	\
+					NULL_TREE);
 #define DEF_FUNCTION_TYPE_VAR_6(ENUM, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
 				ARG6)	\
   builtin_types[(int) ENUM]						\
@@ -1249,6 +1261,7 @@ gfc_init_builtin_functions (void)
 #undef DEF_FUNCTION_TYPE_VAR_0
 #undef DEF_FUNCTION_TYPE_VAR_1
 #undef DEF_FUNCTION_TYPE_VAR_2
+#undef DEF_FUNCTION_TYPE_VAR_5
 #undef DEF_FUNCTION_TYPE_VAR_6
 #undef DEF_FUNCTION_TYPE_VAR_7
 #undef DEF_POINTER_TYPE
diff --git a/gcc/fortran/types.def b/gcc/fortran/types.def
index dd9b8df..9c47785 100644
--- a/gcc/fortran/types.def
+++ b/gcc/fortran/types.def
@@ -283,6 +283,9 @@ DEF_FUNCTION_TYPE_VAR_7 (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_INT_INT_VAR,
 			 BT_VOID, BT_INT, BT_SIZE, BT_PTR, BT_PTR,
 			 BT_PTR, BT_INT, BT_INT)
 
+DEF_FUNCTION_TYPE_VAR_5 (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_VAR,
+			 BT_VOID, BT_INT, BT_SIZE, BT_PTR, BT_PTR, BT_PTR)
+
 DEF_FUNCTION_TYPE_VAR_6 (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_VAR,
 			  BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_SIZE,
 			  BT_PTR, BT_PTR, BT_PTR)
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 4f385b1..8559971 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -10643,6 +10643,14 @@ omp_group_base (omp_mapping_group *grp, unsigned int *chained,
     case GOMP_MAP_ALWAYS_PRESENT_FROM:
     case GOMP_MAP_ALWAYS_PRESENT_TO:
     case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
+    case GOMP_MAP_NONCONTIG_ARRAY_ALLOC:
+    case GOMP_MAP_NONCONTIG_ARRAY_FROM:
+    case GOMP_MAP_NONCONTIG_ARRAY_TO:
+    case GOMP_MAP_NONCONTIG_ARRAY_TOFROM:
+    case GOMP_MAP_NONCONTIG_ARRAY_FORCE_ALLOC:
+    case GOMP_MAP_NONCONTIG_ARRAY_FORCE_FROM:
+    case GOMP_MAP_NONCONTIG_ARRAY_FORCE_TO:
+    case GOMP_MAP_NONCONTIG_ARRAY_FORCE_TOFROM:
     case GOMP_MAP_ALLOC:
     case GOMP_MAP_RELEASE:
     case GOMP_MAP_DELETE:
@@ -12943,6 +12951,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
       || code == OMP_TARGET_DATA
       || code == OMP_TARGET_ENTER_DATA
       || code == OMP_TARGET_EXIT_DATA
+      || code == OMP_TARGET_UPDATE
       || code == OACC_DATA
       || code == OACC_KERNELS
       || code == OACC_PARALLEL
@@ -15052,7 +15061,15 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
 				    : TYPE_SIZE_UNIT (TREE_TYPE (decl));
 	    }
 	  gimplify_omp_ctxp = ctx->outer_context;
-	  if (gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p, NULL,
+	  if (GOMP_MAP_NONCONTIG_ARRAY_P (OMP_CLAUSE_MAP_KIND (c)))
+	    {
+	      gcc_assert (OMP_CLAUSE_SIZE (c)
+			  && TREE_CODE (OMP_CLAUSE_SIZE (c)) == TREE_LIST);
+	      /* For non-contiguous array maps, OMP_CLAUSE_SIZE is a TREE_LIST
+		 of the individual array dimensions, which gimplify_expr doesn't
+		 handle, so skip the call to gimplify_expr here.  */
+	    }
+	  else if (gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p, NULL,
 				  is_gimple_val, fb_rvalue) == GS_ERROR)
 	    {
 	      gimplify_omp_ctxp = ctx;
diff --git a/gcc/omp-builtins.def b/gcc/omp-builtins.def
index f73fb7b..4e7f852 100644
--- a/gcc/omp-builtins.def
+++ b/gcc/omp-builtins.def
@@ -31,7 +31,7 @@ along with GCC; see the file COPYING3.  If not see
    doesn't source those.  */
 
 DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DATA_START, "GOACC_data_start",
-		   BT_FN_VOID_INT_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
+		   BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_VAR, ATTR_NOTHROW_LIST)
 DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DATA_END, "GOACC_data_end",
 		   BT_FN_VOID, ATTR_NOTHROW_LIST)
 DEF_GOACC_BUILTIN (BUILT_IN_GOACC_ENTER_DATA, "GOACC_enter_data",
diff --git a/gcc/omp-expand.cc b/gcc/omp-expand.cc
index 648ede2..2414dd5 100644
--- a/gcc/omp-expand.cc
+++ b/gcc/omp-expand.cc
@@ -10633,6 +10633,19 @@ expand_omp_target (struct omp_region *region)
       gsi_insert_before (&gsi, g, GSI_SAME_STMT);
     }
 
+  /* We assume index >= 3 in gimple_omp_target_data_arg are non-contiguous
+     array descriptor pointer arguments.  */
+  if (t != NULL
+      && TREE_VEC_LENGTH (t) > 3
+      && (start_ix == BUILT_IN_GOACC_DATA_START
+	  || start_ix == BUILT_IN_GOACC_PARALLEL))
+    {
+      gcc_assert ((c = omp_find_clause (clauses, OMP_CLAUSE_MAP))
+		  && GOMP_MAP_NONCONTIG_ARRAY_P (OMP_CLAUSE_MAP_KIND (c)));
+      for (int i = 3; i < TREE_VEC_LENGTH (t); i++)
+	args.safe_push (TREE_VEC_ELT (t, i));
+    }
+
   g = gimple_build_call_vec (builtin_decl_explicit (start_ix), args);
   gimple_set_location (g, gimple_location (entry_stmt));
   gsi_insert_before (&gsi, g, GSI_SAME_STMT);
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index e1036ad..ca1767b 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -968,6 +968,123 @@ omp_copy_decl (tree var, copy_body_data *cb)
   return error_mark_node;
 }
 
+/* Helper function for create_noncontig_array_descr_type(), to append a new field
+   to a record type.  */
+
+static void
+append_field_to_record_type (tree record_type, tree fld_ident, tree fld_type)
+{
+  tree *p, fld = build_decl (UNKNOWN_LOCATION, FIELD_DECL, fld_ident, fld_type);
+  DECL_CONTEXT (fld) = record_type;
+
+  for (p = &TYPE_FIELDS (record_type); *p; p = &DECL_CHAIN (*p))
+    ;
+  *p = fld;
+}
+
+/* Create type for non-contiguous array descriptor. Returns created type, and
+   returns the number of dimensions in *DIM_NUM.  */
+
+static tree
+create_noncontig_array_descr_type (tree dims, int *dim_num)
+{
+  int n = 0;
+  tree array_descr_type, name, x;
+  gcc_assert (TREE_CODE (dims) == TREE_LIST);
+
+  array_descr_type = lang_hooks.types.make_type (RECORD_TYPE);
+  name = create_tmp_var_name (".omp_noncontig_array_descr_type");
+  name = build_decl (UNKNOWN_LOCATION, TYPE_DECL, name, array_descr_type);
+  DECL_ARTIFICIAL (name) = 1;
+  DECL_NAMELESS (name) = 1;
+  TYPE_NAME (array_descr_type) = name;
+  TYPE_ARTIFICIAL (array_descr_type) = 1;
+
+  /* Number of dimensions.  */
+  append_field_to_record_type (array_descr_type, get_identifier ("__dim_num"),
+			       sizetype);
+
+  for (x = dims; x; x = TREE_CHAIN (x), n++)
+    {
+      char *fldname;
+      /* One for the start index.  */
+      ASM_FORMAT_PRIVATE_NAME (fldname, "__dim_base", n);
+      append_field_to_record_type (array_descr_type, get_identifier (fldname),
+				   sizetype);
+      /* One for the length.  */
+      ASM_FORMAT_PRIVATE_NAME (fldname, "__dim_length", n);
+      append_field_to_record_type (array_descr_type, get_identifier (fldname),
+				   sizetype);
+      /* One for the element size.  */
+      ASM_FORMAT_PRIVATE_NAME (fldname, "__dim_elem_size", n);
+      append_field_to_record_type (array_descr_type, get_identifier (fldname),
+				   sizetype);
+      /* One for is_array flag.  */
+      ASM_FORMAT_PRIVATE_NAME (fldname, "__dim_is_array", n);
+      append_field_to_record_type (array_descr_type, get_identifier (fldname),
+				   sizetype);
+    }
+
+  layout_type (array_descr_type);
+  *dim_num = n;
+  return array_descr_type;
+}
+
+/* Generate code sequence for initializing non-contiguous array descriptor.  */
+
+static void
+create_noncontig_array_descr_init_code (tree array_descr, tree array_var,
+					tree dimensions, int dim_num,
+					gimple_seq *ilist)
+{
+  tree fld, fldref;
+  tree array_descr_type = TREE_TYPE (array_descr);
+  tree dim_type = TREE_TYPE (array_var);
+
+  if (TREE_CODE (dim_type) == REFERENCE_TYPE)
+    dim_type = TREE_TYPE (dim_type);
+
+  fld = TYPE_FIELDS (array_descr_type);
+  fldref = omp_build_component_ref (array_descr, fld);
+  gimplify_assign (fldref, build_int_cst (sizetype, dim_num), ilist);
+
+  while (dimensions)
+    {
+      tree dim_base = fold_convert (sizetype, TREE_PURPOSE (dimensions));
+      tree dim_length = fold_convert (sizetype, TREE_VALUE (dimensions));
+      tree dim_elem_size = TYPE_SIZE_UNIT (TREE_TYPE (dim_type));
+      tree dim_is_array = (TREE_CODE (dim_type) == ARRAY_TYPE
+			   ? integer_one_node : integer_zero_node);
+      /* Set base.  */
+      fld = TREE_CHAIN (fld);
+      fldref = omp_build_component_ref (array_descr, fld);
+      dim_base = fold_build2 (MULT_EXPR, sizetype, dim_base, dim_elem_size);
+      gimplify_assign (fldref, dim_base, ilist);
+
+      /* Set length.  */
+      fld = TREE_CHAIN (fld);
+      fldref = omp_build_component_ref (array_descr, fld);
+      dim_length = fold_build2 (MULT_EXPR, sizetype, dim_length, dim_elem_size);
+      gimplify_assign (fldref, dim_length, ilist);
+
+      /* Set elem_size.  */
+      fld = TREE_CHAIN (fld);
+      fldref = omp_build_component_ref (array_descr, fld);
+      dim_elem_size = fold_convert (sizetype, dim_elem_size);
+      gimplify_assign (fldref, dim_elem_size, ilist);
+
+      /* Set is_array flag.  */
+      fld = TREE_CHAIN (fld);
+      fldref = omp_build_component_ref (array_descr, fld);
+      dim_is_array = fold_convert (sizetype, dim_is_array);
+      gimplify_assign (fldref, dim_is_array, ilist);
+
+      dimensions = TREE_CHAIN (dimensions);
+      dim_type = TREE_TYPE (dim_type);
+    }
+  gcc_assert (TREE_CHAIN (fld) == NULL_TREE);
+}
+
 /* Create a new context, with OUTER_CTX being the surrounding context.  */
 
 static omp_context *
@@ -1670,6 +1787,38 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	      install_var_local (decl, ctx);
 	      break;
 	    }
+
+	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+	      && GOMP_MAP_NONCONTIG_ARRAY_P (OMP_CLAUSE_MAP_KIND (c)))
+	    {
+	      tree array_decl = OMP_CLAUSE_DECL (c);
+	      tree array_type = TREE_TYPE (array_decl);
+	      bool by_ref = (TREE_CODE (array_type) == ARRAY_TYPE
+			     ? true : false);
+
+	      /* Checking code to ensure we only have arrays at top dimension.
+		 This limitation might be lifted in the future. See PR76639.  */
+	      if (TREE_CODE (array_type) == REFERENCE_TYPE)
+		array_type = TREE_TYPE (array_type);
+	      tree t = array_type, prev_t = NULL_TREE;
+	      while (t)
+		{
+		  if (TREE_CODE (t) == ARRAY_TYPE && prev_t)
+		    {
+		      error_at (gimple_location (ctx->stmt), "array types are"
+				" only allowed at outermost dimension of"
+				" non-contiguous array");
+		      break;
+		    }
+		  prev_t = t;
+		  t = TREE_TYPE (t);
+		}
+
+	      install_var_field (array_decl, by_ref, 3, ctx);
+	      install_var_local (array_decl, ctx);
+	      break;
+	    }
+
 	  if (DECL_P (decl))
 	    {
 	      if (DECL_SIZE (decl)
@@ -3119,6 +3268,50 @@ scan_omp_single (gomp_single *stmt, omp_context *outer_ctx)
     layout_type (ctx->record_type);
 }
 
+/* Reorder clauses so that non-contiguous array map clauses are placed at the very
+   front of the chain.  */
+
+static void
+reorder_noncontig_array_clauses (tree *clauses_ptr)
+{
+  tree c, clauses = *clauses_ptr;
+  tree prev_clause = NULL_TREE, next_clause;
+  tree array_clauses = NULL_TREE, array_clauses_tail = NULL_TREE;
+
+  for (c = clauses; c; c = next_clause)
+    {
+      next_clause = OMP_CLAUSE_CHAIN (c);
+
+      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+	  && GOMP_MAP_NONCONTIG_ARRAY_P (OMP_CLAUSE_MAP_KIND (c)))
+	{
+	  /* Unchain c from clauses.  */
+	  if (c == clauses)
+	    clauses = next_clause;
+
+	  /* Link on to array_clauses.  */
+	  if (array_clauses_tail)
+	    OMP_CLAUSE_CHAIN (array_clauses_tail) = c;
+	  else
+	    array_clauses = c;
+	  array_clauses_tail = c;
+
+	  if (prev_clause)
+	    OMP_CLAUSE_CHAIN (prev_clause) = next_clause;
+	  continue;
+	}
+
+      prev_clause = c;
+    }
+
+  /* Place non-contiguous array clauses at the start of the clause list.  */
+  if (array_clauses)
+    {
+      OMP_CLAUSE_CHAIN (array_clauses_tail) = clauses;
+      *clauses_ptr = array_clauses;
+    }
+}
+
 /* Scan a GIMPLE_OMP_TARGET.  */
 
 static void
@@ -3127,7 +3320,6 @@ scan_omp_target (gomp_target *stmt, omp_context *outer_ctx)
   omp_context *ctx;
   tree name;
   bool offloaded = is_gimple_omp_offloaded (stmt);
-  tree clauses = gimple_omp_target_clauses (stmt);
 
   ctx = new_omp_context (stmt, outer_ctx);
   ctx->field_map = splay_tree_new (splay_tree_compare_pointers, 0, 0);
@@ -3140,6 +3332,14 @@ scan_omp_target (gomp_target *stmt, omp_context *outer_ctx)
   TYPE_NAME (ctx->record_type) = name;
   TYPE_ARTIFICIAL (ctx->record_type) = 1;
 
+  /* If is OpenACC construct, put non-contiguous array clauses (if any)
+     in front of clause chain. The runtime can then test the first to see
+     if the additional map processing for them is required.  */
+  if (is_gimple_omp_oacc (stmt))
+    reorder_noncontig_array_clauses (gimple_omp_target_clauses_ptr (stmt));
+
+  tree clauses = gimple_omp_target_clauses (stmt);
+
   if (offloaded)
     {
       create_omp_child_function (ctx, false);
@@ -12797,6 +12997,15 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	  case GOMP_MAP_FORCE_TOFROM:
 	  case GOMP_MAP_FORCE_DEVICEPTR:
 	  case GOMP_MAP_DEVICE_RESIDENT:
+	  case GOMP_MAP_NONCONTIG_ARRAY_TO:
+	  case GOMP_MAP_NONCONTIG_ARRAY_FROM:
+	  case GOMP_MAP_NONCONTIG_ARRAY_TOFROM:
+	  case GOMP_MAP_NONCONTIG_ARRAY_FORCE_TO:
+	  case GOMP_MAP_NONCONTIG_ARRAY_FORCE_FROM:
+	  case GOMP_MAP_NONCONTIG_ARRAY_FORCE_TOFROM:
+	  case GOMP_MAP_NONCONTIG_ARRAY_ALLOC:
+	  case GOMP_MAP_NONCONTIG_ARRAY_FORCE_ALLOC:
+	  case GOMP_MAP_NONCONTIG_ARRAY_FORCE_PRESENT:
 	  case GOMP_MAP_LINK:
 	  case GOMP_MAP_FORCE_DETACH:
 	    gcc_assert (is_gimple_omp_oacc (stmt));
@@ -12881,8 +13090,15 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 			   && is_gimple_omp_oacc (ctx->stmt)
 			   && OMP_CLAUSE_MAP_IN_REDUCTION (c)))
 	  {
-	    x = build_receiver_ref (var, true, ctx);
+	    tree var_type = TREE_TYPE (var);
 	    tree new_var = lookup_decl (var, ctx);
+	    bool rcv_by_ref =
+	      (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+	       && GOMP_MAP_NONCONTIG_ARRAY_P (OMP_CLAUSE_MAP_KIND (c))
+	       && TREE_CODE (var_type) != ARRAY_TYPE
+	       ? false : true);
+
+	    x = build_receiver_ref (var, rcv_by_ref, ctx);
 
 	    if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 		&& OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
@@ -13151,6 +13367,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
       vec_alloc (vkind, map_cnt);
       unsigned int map_idx = 0;
 
+      vec<tree> nca_descrs = vNULL;
+
       for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
 	switch (OMP_CLAUSE_CODE (c))
 	  {
@@ -13315,6 +13533,29 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 		    avar = build_fold_addr_expr (avar);
 		    gimplify_assign (x, avar, &ilist);
 		  }
+		else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+			 && GOMP_MAP_NONCONTIG_ARRAY_P (OMP_CLAUSE_MAP_KIND (c)))
+		  {
+		    int dim_num;
+		    tree dimensions = OMP_CLAUSE_SIZE (c);
+
+		    tree array_descr_type =
+		      create_noncontig_array_descr_type (dimensions, &dim_num);
+		    tree array_descr =
+		      create_tmp_var_raw (array_descr_type,
+					  ".omp_noncontig_array_descr");
+		    TREE_ADDRESSABLE (array_descr) = 1;
+		    TREE_STATIC (array_descr) = 1;
+		    gimple_add_tmp_var (array_descr);
+
+		    create_noncontig_array_descr_init_code
+		      (array_descr, ovar, dimensions, dim_num, &ilist);
+		    nca_descrs.safe_push (build_fold_addr_expr (array_descr));
+
+		    gimplify_assign (x, (TREE_CODE (TREE_TYPE (ovar)) == ARRAY_TYPE
+					 ? build_fold_addr_expr (ovar) : ovar),
+				     &ilist);
+		  }
 		else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE)
 		  {
 		    gcc_assert (is_gimple_omp_oacc (ctx->stmt));
@@ -13387,6 +13628,9 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 		  s = TREE_TYPE (s);
 		s = TYPE_SIZE_UNIT (s);
 	      }
+	    else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+		     && GOMP_MAP_NONCONTIG_ARRAY_P (OMP_CLAUSE_MAP_KIND (c)))
+	      s = NULL_TREE;
 	    else
 	      s = OMP_CLAUSE_SIZE (c);
 	    if (s == NULL_TREE)
@@ -13750,6 +13994,19 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 
       gcc_assert (map_idx == map_cnt);
 
+      unsigned nca_num = nca_descrs.length ();
+      if (nca_num > 0)
+	{
+	  tree nca, t = gimple_omp_target_data_arg (stmt);
+	  int i, oldlen = TREE_VEC_LENGTH (t);
+	  tree nt = make_tree_vec (oldlen + nca_num);
+	  for (i = 0; i < oldlen; i++)
+	    TREE_VEC_ELT (nt, i) = TREE_VEC_ELT (t, i);
+	  for (i = 0; nca_descrs.iterate (i, &nca); i++)
+	    TREE_VEC_ELT (nt, oldlen + i) = nca;
+	  gimple_omp_target_set_data_arg (stmt, nt);
+	}
+
       if (!deep_map_cnt)
 	{
 	  DECL_INITIAL (TREE_VEC_ELT (t, 1))
diff --git a/gcc/testsuite/c-c++-common/goacc/data-clause-1.c b/gcc/testsuite/c-c++-common/goacc/data-clause-1.c
index 9952ac4..b78f691 100644
--- a/gcc/testsuite/c-c++-common/goacc/data-clause-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/data-clause-1.c
@@ -98,7 +98,7 @@ foo (int g[3][10], int h[4][8], int i[2][10], int j[][9],
     bar (&j2[0][0]);
   #pragma acc parallel copy(q[1:2])
     ;
-  #pragma acc parallel copy(q[3:5][:10]) /* { dg-error "array section is not contiguous" } */
+  #pragma acc parallel copy(q[3:5][:10])
     ;
   #pragma acc parallel copy(r[3:][2:1][1:2])
     ;
diff --git a/gcc/testsuite/c-c++-common/goacc/noncontig_array-1.c b/gcc/testsuite/c-c++-common/goacc/noncontig_array-1.c
new file mode 100644
index 0000000..fe7480a
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/noncontig_array-1.c
@@ -0,0 +1,26 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+void foo (void)
+{
+  int array_of_array[10][10];
+  int **ptr_to_ptr;
+  int *array_of_ptr[10];
+  int (*ptr_to_array)[10];
+ 
+  #pragma acc parallel copy (array_of_array[2:4][0:10])
+    array_of_array[5][5] = 1;
+
+  #pragma acc parallel copy (ptr_to_ptr[2:4][1:7])
+    ptr_to_ptr[5][5] = 1;
+
+  #pragma acc parallel copy (array_of_ptr[2:4][1:7])
+    array_of_ptr[5][5] = 1;
+
+  #pragma acc parallel copy (ptr_to_array[2:4][1:7]) /* { dg-error "array section is not contiguous in 'map' clause" } */
+    ptr_to_array[5][5] = 1;
+}
+/* { dg-final { scan-tree-dump-times {#pragma omp target oacc_parallel map\(tofrom:array_of_array} 1 gimple } } */
+/* { dg-final { scan-tree-dump-times {#pragma omp target oacc_parallel map\(tofrom,noncontig_array:ptr_to_ptr \[dimensions: 2 4, 1 7\]} 1 gimple } } */
+/* { dg-final { scan-tree-dump-times {#pragma omp target oacc_parallel map\(tofrom,noncontig_array:array_of_ptr \[dimensions: 2 4, 1 7\]} 1 gimple } } */
+/* { dg-final { scan-tree-dump-times {#pragma omp target oacc_parallel map\(tofrom,noncontig_array:ptr_to_array \[dimensions: 2 4, 1 7\]} 1 gimple { xfail *-*-* } } } */
diff --git a/gcc/testsuite/g++.dg/goacc/data-clause-1.C b/gcc/testsuite/g++.dg/goacc/data-clause-1.C
index 07ef6ae..daea3f4 100644
--- a/gcc/testsuite/g++.dg/goacc/data-clause-1.C
+++ b/gcc/testsuite/g++.dg/goacc/data-clause-1.C
@@ -99,7 +99,7 @@ foo (int g[3][10], int h[4][8], int i[2][10], int j[][9],
     bar (&j2[0][0]);
   #pragma acc parallel copy(q[1:2])
     ;
-  #pragma acc parallel copy(q[3:5][:10]) /* { dg-error "array section is not contiguous" } */
+  #pragma acc parallel copy(q[3:5][:10])
     ;
   #pragma acc parallel copy(r[3:][2:1][1:2])
     ;
diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc
index c1a21e7..662236f 100644
--- a/gcc/tree-pretty-print.cc
+++ b/gcc/tree-pretty-print.cc
@@ -1122,6 +1122,33 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
 	case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
 	  pp_string (pp, "always,present,tofrom");
 	  break;
+	case GOMP_MAP_NONCONTIG_ARRAY_TO:
+	  pp_string (pp, "to,noncontig_array");
+	  break;
+	case GOMP_MAP_NONCONTIG_ARRAY_FROM:
+	  pp_string (pp, "from,noncontig_array");
+	  break;
+	case GOMP_MAP_NONCONTIG_ARRAY_TOFROM:
+	  pp_string (pp, "tofrom,noncontig_array");
+	  break;
+	case GOMP_MAP_NONCONTIG_ARRAY_FORCE_TO:
+	  pp_string (pp, "force_to,noncontig_array");
+	  break;
+	case GOMP_MAP_NONCONTIG_ARRAY_FORCE_FROM:
+	  pp_string (pp, "force_from,noncontig_array");
+	  break;
+	case GOMP_MAP_NONCONTIG_ARRAY_FORCE_TOFROM:
+	  pp_string (pp, "force_tofrom,noncontig_array");
+	  break;
+	case GOMP_MAP_NONCONTIG_ARRAY_ALLOC:
+	  pp_string (pp, "alloc,noncontig_array");
+	  break;
+	case GOMP_MAP_NONCONTIG_ARRAY_FORCE_ALLOC:
+	  pp_string (pp, "force_alloc,noncontig_array");
+	  break;
+	case GOMP_MAP_NONCONTIG_ARRAY_FORCE_PRESENT:
+	  pp_string (pp, "force_present,noncontig_array");
+	  break;
 	default:
 	  gcc_unreachable ();
 	}
@@ -1132,8 +1159,15 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
       if (OMP_CLAUSE_SIZE (clause))
 	{
 	  switch (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP
-		  ? OMP_CLAUSE_MAP_KIND (clause) : GOMP_MAP_TO)
+		  ? (GOMP_MAP_NONCONTIG_ARRAY_P (OMP_CLAUSE_MAP_KIND (clause))
+		     ? GOMP_MAP_NONCONTIG_ARRAY
+		     : OMP_CLAUSE_MAP_KIND (clause))
+		  : GOMP_MAP_TO)
 	    {
+	    case GOMP_MAP_NONCONTIG_ARRAY:
+	      gcc_assert (TREE_CODE (OMP_CLAUSE_SIZE (clause)) == TREE_LIST);
+	      pp_string (pp, " [dimensions: ");
+	      break;
 	    case GOMP_MAP_POINTER:
 	    case GOMP_MAP_FIRSTPRIVATE_POINTER:
 	    case GOMP_MAP_FIRSTPRIVATE_REFERENCE: