[OpenACC,v3] Non-contiguous array support for OpenACC data clauses

Message ID 6b17767f-9b7d-50a7-c52f-e80bf1be991e@mentor.com
State New
Headers show
Series
  • [OpenACC,v3] Non-contiguous array support for OpenACC data clauses
Related show

Commit Message

Chung-Lin Tang Nov. 26, 2019, 2:49 p.m.
Hi Thomas,
this is a reorg of the last non-contiguous arrays patch. You'll notice that:

(1) A large part of the code has been pulled into oacc-parallel.c, with most
of the data structure declarations in oacc-int.h.

(2) target.c only contains relatively little code from gomp_map_vars_internal
that processes what GOACC_parallel_keyed/data_start gives it.

(3) Instead of directly passed in the map pointer, the array descriptor
pointers are now passed to GOACC_parallel_keyed/data_start using varargs.
(I believe the adding of '...' to GOACC_data_start does not break any
compatiblity)

(4) Along the way, I've added a 'gomp_map_vars_openacc' for specializing our
uses, which should shave off quite some code through inlining.

The GOMP_MAP_NONCONTIG_ARRAY_P maps are still placed at the beginning of the
recieved map sequence in this patch. It should still be relatively easy to
use a GOACC_FLAG_* to do so if deemed better before committing.

Thanks,
Chung-Lin

         PR other/76739

         gcc/c/
         * c-typeck.c (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.
         (handle_omp_array_sections): Adjust handle_omp_array_sections_1 call,
         handle non-contiguous case to create dynamic array map.

         gcc/cp/
         * semantics.c (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.
         (handle_omp_array_sections): Adjust handle_omp_array_sections_1 call,
         handle non-contiguous case to create dynamic array map.

         gcc/fortran/
         * f95-lang.c (DEF_FUNCTION_TYPE_VAR_5): New symbol.
         * types.def (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_VAR): New type.

         gcc/
         * builtin-types.def (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_VAR): New type.
         * omp-builtins.def (BUILT_IN_GOACC_DATA_START): Adjust function type
         to new BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_VAR.
         * gimplify.c (gimplify_scan_omp_clauses): Skip gimplification of
         OMP_CLAUSE_SIZE of non-contiguous array maps (which is a TREE_LIST).
         * omp-expand.c (expand_omp_target): Add non-contiguous array descriptor
         pointers to variadic arguments.
         * omp-low.c (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.

         gcc/testsuite/
	* c-c++-common/goacc/noncontig_array-1.c: New test.

	libgomp/
         * libgomp_g.h (GOACC_data_start): Add variadic '...' to declaration.
         * libgomp.h (gomp_map_vars_openacc): New function 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.

Patch

Index: include/gomp-constants.h
===================================================================
--- include/gomp-constants.h	(revision 278656)
+++ include/gomp-constants.h	(working copy)
@@ -40,6 +40,7 @@ 
 #define GOMP_MAP_FLAG_SPECIAL_0		(1 << 2)
 #define GOMP_MAP_FLAG_SPECIAL_1		(1 << 3)
 #define GOMP_MAP_FLAG_SPECIAL_2		(1 << 4)
+#define GOMP_MAP_FLAG_SPECIAL_3		(1 << 5)
 #define GOMP_MAP_FLAG_SPECIAL		(GOMP_MAP_FLAG_SPECIAL_1 \
 					 | GOMP_MAP_FLAG_SPECIAL_0)
 /* Flag to force a specific behavior (or else, trigger a run-time error).  */
@@ -127,6 +128,26 @@  enum gomp_map_kind
     /* Decrement usage count and deallocate if zero.  */
     GOMP_MAP_RELEASE =			(GOMP_MAP_FLAG_SPECIAL_2
 					 | GOMP_MAP_DELETE),
+    /* Mapping kinds for non-contiguous arrays.  */
+    GOMP_MAP_NONCONTIG_ARRAY =		(GOMP_MAP_FLAG_SPECIAL_3),
+    GOMP_MAP_NONCONTIG_ARRAY_TO =	(GOMP_MAP_NONCONTIG_ARRAY
+					 | GOMP_MAP_TO),
+    GOMP_MAP_NONCONTIG_ARRAY_FROM =	(GOMP_MAP_NONCONTIG_ARRAY
+					 | GOMP_MAP_FROM),
+    GOMP_MAP_NONCONTIG_ARRAY_TOFROM =	(GOMP_MAP_NONCONTIG_ARRAY
+					 | GOMP_MAP_TOFROM),
+    GOMP_MAP_NONCONTIG_ARRAY_FORCE_TO =	(GOMP_MAP_NONCONTIG_ARRAY_TO
+					 | GOMP_MAP_FLAG_FORCE),
+    GOMP_MAP_NONCONTIG_ARRAY_FORCE_FROM =	(GOMP_MAP_NONCONTIG_ARRAY_FROM
+						 | GOMP_MAP_FLAG_FORCE),
+    GOMP_MAP_NONCONTIG_ARRAY_FORCE_TOFROM =	(GOMP_MAP_NONCONTIG_ARRAY_TOFROM
+						 | GOMP_MAP_FLAG_FORCE),
+    GOMP_MAP_NONCONTIG_ARRAY_ALLOC =		(GOMP_MAP_NONCONTIG_ARRAY
+						 | GOMP_MAP_ALLOC),
+    GOMP_MAP_NONCONTIG_ARRAY_FORCE_ALLOC =	(GOMP_MAP_NONCONTIG_ARRAY
+						 | GOMP_MAP_FORCE_ALLOC),
+    GOMP_MAP_NONCONTIG_ARRAY_FORCE_PRESENT =	(GOMP_MAP_NONCONTIG_ARRAY
+						 | GOMP_MAP_FORCE_PRESENT),
 
     /* Internal to GCC, not used in libgomp.  */
     /* Do not map, but pointer assign a pointer instead.  */
@@ -155,6 +176,8 @@  enum gomp_map_kind
 #define GOMP_MAP_ALWAYS_P(X) \
   (GOMP_MAP_ALWAYS_TO_P (X) || ((X) == GOMP_MAP_ALWAYS_FROM))
 
+#define GOMP_MAP_NONCONTIG_ARRAY_P(X) \
+  ((X) & GOMP_MAP_NONCONTIG_ARRAY)
 
 /* Asynchronous behavior.  Keep in sync with
    libgomp/{openacc.h,openacc.f90,openacc_lib.h}:acc_async_t.  */
Index: gcc/builtin-types.def
===================================================================
--- gcc/builtin-types.def	(revision 278656)
+++ gcc/builtin-types.def	(working copy)
@@ -822,6 +822,9 @@  DEF_FUNCTION_TYPE_VAR_5 (BT_FN_INT_STRING_SIZE_INT
 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)
Index: gcc/c/c-typeck.c
===================================================================
--- gcc/c/c-typeck.c	(revision 278656)
+++ gcc/c/c-typeck.c	(working copy)
@@ -12876,12 +12876,14 @@  c_finish_omp_cancellation_point (location_t loc, t
    <= 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;
   if (TREE_CODE (t) != TREE_LIST)
@@ -12966,7 +12968,8 @@  handle_omp_array_sections_1 (tree c, tree t, vec<t
     }
 
   ret = handle_omp_array_sections_1 (c, TREE_CHAIN (t), 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;
 
@@ -13173,14 +13176,21 @@  handle_omp_array_sections_1 (tree c, tree t, vec<t
 	  return error_mark_node;
 	}
       /* If there is a pointer type anywhere but in the very first
-	 array-section-subscript, the array section can't be contiguous.  */
+	 array-section-subscript, the array section can't be contiguous.
+	 Note that OpenACC does accept these kinds of non-contiguous pointer
+	 based arrays.  */
       if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND
 	  && TREE_CODE (TREE_CHAIN (t)) == TREE_LIST)
 	{
-	  error_at (OMP_CLAUSE_LOCATION (c),
-		    "array section is not contiguous in %qs clause",
-		    omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
-	  return error_mark_node;
+	  if (ort == C_ORT_ACC)
+	    non_contiguous = true;
+	  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;
+	    }
 	}
     }
   else
@@ -13209,6 +13219,7 @@  handle_omp_array_sections (tree c, enum c_omp_regi
 {
   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
@@ -13218,7 +13229,7 @@  handle_omp_array_sections (tree c, enum c_omp_regi
     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)
@@ -13251,6 +13262,7 @@  handle_omp_array_sections (tree c, enum c_omp_regi
       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;
@@ -13274,6 +13286,13 @@  handle_omp_array_sections (tree c, enum c_omp_regi
 	    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))
@@ -13370,6 +13389,14 @@  handle_omp_array_sections (tree c, enum c_omp_regi
 		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
Index: gcc/cp/semantics.c
===================================================================
--- gcc/cp/semantics.c	(revision 278656)
+++ gcc/cp/semantics.c	(working copy)
@@ -4735,12 +4735,14 @@  omp_privatize_field (tree t, bool shared)
    <= 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;
   if (TREE_CODE (t) != TREE_LIST)
@@ -4825,7 +4827,8 @@  handle_omp_array_sections_1 (tree c, tree t, vec<t
       && TREE_CODE (TREE_CHAIN (t)) == FIELD_DECL)
     TREE_CHAIN (t) = omp_privatize_field (TREE_CHAIN (t), false);
   ret = handle_omp_array_sections_1 (c, TREE_CHAIN (t), 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;
 
@@ -5044,14 +5047,21 @@  handle_omp_array_sections_1 (tree c, tree t, vec<t
 	  return error_mark_node;
 	}
       /* If there is a pointer type anywhere but in the very first
-	 array-section-subscript, the array section can't be contiguous.  */
+	 array-section-subscript, the array section can't be contiguous.
+	 Note that OpenACC does accept these kinds of non-contiguous pointer
+	 based arrays.  */
       if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND
 	  && TREE_CODE (TREE_CHAIN (t)) == TREE_LIST)
 	{
-	  error_at (OMP_CLAUSE_LOCATION (c),
-		    "array section is not contiguous in %qs clause",
-		    omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
-	  return error_mark_node;
+	  if (ort == C_ORT_ACC)
+	    non_contiguous = true;
+	  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;
+	    }
 	}
     }
   else
@@ -5091,6 +5101,7 @@  handle_omp_array_sections (tree c, enum c_omp_regi
 {
   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
@@ -5100,7 +5111,7 @@  handle_omp_array_sections (tree c, enum c_omp_regi
     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)
@@ -5134,6 +5145,7 @@  handle_omp_array_sections (tree c, enum c_omp_regi
       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;
@@ -5159,6 +5171,13 @@  handle_omp_array_sections (tree c, enum c_omp_regi
 	    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))
@@ -5250,6 +5269,14 @@  handle_omp_array_sections (tree c, enum c_omp_regi
 	}
       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
Index: gcc/fortran/f95-lang.c
===================================================================
--- gcc/fortran/f95-lang.c	(revision 278656)
+++ gcc/fortran/f95-lang.c	(working copy)
@@ -632,6 +632,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, \
Index: gcc/fortran/types.def
===================================================================
--- gcc/fortran/types.def	(revision 278656)
+++ gcc/fortran/types.def	(working copy)
@@ -269,6 +269,9 @@  DEF_FUNCTION_TYPE_VAR_7 (BT_FN_VOID_INT_SIZE_PTR_P
 			 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)
Index: gcc/gimplify.c
===================================================================
--- gcc/gimplify.c	(revision 278656)
+++ gcc/gimplify.c	(working copy)
@@ -8641,9 +8641,17 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_se
 	  if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
 	    OMP_CLAUSE_SIZE (c) = DECL_P (decl) ? DECL_SIZE_UNIT (decl)
 				  : TYPE_SIZE_UNIT (TREE_TYPE (decl));
-	  if (gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p,
-			     NULL, is_gimple_val, fb_rvalue) == GS_ERROR)
+	  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)
+	    {
 	      remove = true;
 	      break;
 	    }
Index: gcc/omp-builtins.def
===================================================================
--- gcc/omp-builtins.def	(revision 278656)
+++ gcc/omp-builtins.def	(working copy)
@@ -32,7 +32,7 @@  along with GCC; see the file COPYING3.  If not see
 DEF_GOACC_BUILTIN (BUILT_IN_ACC_GET_DEVICE_TYPE, "acc_get_device_type",
 		   BT_FN_INT, ATTR_NOTHROW_LIST)
 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_EXIT_DATA, "GOACC_enter_exit_data",
Index: gcc/omp-expand.c
===================================================================
--- gcc/omp-expand.c	(revision 278656)
+++ gcc/omp-expand.c	(working copy)
@@ -8458,6 +8458,21 @@  expand_omp_target (struct omp_region *region)
   if (tagging)
     /*  Push terminal marker - zero.  */
     args.safe_push (oacc_launch_pack (0, NULL_TREE, 0));
+  
+  /* We assume index >= 3 in gimple_omp_target_data_arg are non-contiguous
+     array descriptor pointer arguments.  */
+  if (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)));
+      //if (tagging)
+      //args.safe_push (oacc_launch_pack (GOMP_LAUNCH_NONCONTIG_ARRAYS,
+      //				  NULL_TREE, 0));
+      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));
Index: gcc/omp-low.c
===================================================================
--- gcc/omp-low.c	(revision 278656)
+++ gcc/omp-low.c	(working copy)
@@ -903,6 +903,141 @@  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 decl, 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;
+
+  /* Main starting pointer/array.  */
+  /*
+  tree main_var_type = TREE_TYPE (decl);
+  if (TREE_CODE (main_var_type) == REFERENCE_TYPE)
+    main_var_type = TREE_TYPE (main_var_type);
+  append_field_to_record_type (array_descr_type, DECL_NAME (decl),
+			       (TREE_CODE (TREE_TYPE (decl)) == POINTER_TYPE
+				? main_var_type
+				: build_pointer_type (main_var_type)));
+  */
+  
+  /* 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);
+
+  fld = TYPE_FIELDS (array_descr_type);
+  /*
+  fldref = omp_build_component_ref (array_descr, fld);
+  gimplify_assign (fldref, (TREE_CODE (dim_type) == ARRAY_TYPE
+			    ? build_fold_addr_expr (array_var) : array_var),
+		   ilist);
+  */
+  if (TREE_CODE (dim_type) == REFERENCE_TYPE)
+    dim_type = TREE_TYPE (dim_type);
+
+  //fld = TREE_CHAIN (fld);
+  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 *
@@ -1385,6 +1520,38 @@  scan_sharing_clauses (tree clauses, omp_context *c
 	      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)
@@ -2697,6 +2864,50 @@  scan_omp_single (gomp_single *stmt, omp_context *o
     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
@@ -2705,7 +2916,6 @@  scan_omp_target (gomp_target *stmt, omp_context *o
   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);
@@ -2724,6 +2934,14 @@  scan_omp_target (gomp_target *stmt, omp_context *o
       gimple_omp_target_set_child_fn (stmt, ctx->cb.dst_fn);
     }
 
+  /* 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);
+  
   scan_sharing_clauses (clauses, ctx);
   scan_omp (gimple_omp_body_ptr (stmt), ctx);
 
@@ -11438,6 +11656,15 @@  lower_omp_target (gimple_stmt_iterator *gsi_p, omp
 	  case GOMP_MAP_FORCE_PRESENT:
 	  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:
 	    gcc_assert (is_gimple_omp_oacc (stmt));
 	    break;
@@ -11500,7 +11727,14 @@  lower_omp_target (gimple_stmt_iterator *gsi_p, omp
 	if (offloaded && !(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 			   && OMP_CLAUSE_MAP_IN_REDUCTION (c)))
 	  {
-	    x = build_receiver_ref (var, true, ctx);
+	    tree var_type = TREE_TYPE (var);
+	    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);
 	    tree new_var = lookup_decl (var, ctx);
 
 	    if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
@@ -11674,6 +11908,8 @@  lower_omp_target (gimple_stmt_iterator *gsi_p, omp
       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))
 	  {
@@ -11750,6 +11986,28 @@  lower_omp_target (gimple_stmt_iterator *gsi_p, omp
 		    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 (OMP_CLAUSE_DECL (c),
+							 dimensions, &dim_num);
+		    tree array_descr =
+		      create_tmp_var_raw (array_descr_type,
+					  ".omp_noncontig_array_descr");
+		    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));
@@ -11821,6 +12079,9 @@  lower_omp_target (gimple_stmt_iterator *gsi_p, omp
 		  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)
@@ -12072,6 +12333,19 @@  lower_omp_target (gimple_stmt_iterator *gsi_p, omp
 
       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);
+	}
+
       DECL_INITIAL (TREE_VEC_ELT (t, 1))
 	= build_constructor (TREE_TYPE (TREE_VEC_ELT (t, 1)), vsize);
       DECL_INITIAL (TREE_VEC_ELT (t, 2))
Index: gcc/testsuite/c-c++-common/goacc/noncontig_array-1.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc/noncontig_array-1.c	(nonexistent)
+++ gcc/testsuite/c-c++-common/goacc/noncontig_array-1.c	(working copy)
@@ -0,0 +1,25 @@ 
+/* { dg-do compile } */
+
+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 *-*-* } } } */
Index: gcc/tree-pretty-print.c
===================================================================
--- gcc/tree-pretty-print.c	(revision 278656)
+++ gcc/tree-pretty-print.c	(working copy)
@@ -849,6 +849,33 @@  dump_omp_clause (pretty_printer *pp, tree clause,
 	case GOMP_MAP_LINK:
 	  pp_string (pp, "link");
 	  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 ();
 	}
@@ -859,8 +886,15 @@  dump_omp_clause (pretty_printer *pp, tree clause,
       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:
Index: libgomp/libgomp.h
===================================================================
--- libgomp/libgomp.h	(revision 278656)
+++ libgomp/libgomp.h	(working copy)
@@ -1167,6 +1167,10 @@  extern struct target_mem_desc *gomp_map_vars_async
 						    size_t, void **, void **,
 						    size_t *, void *, bool,
 						    enum gomp_map_vars_kind);
+extern struct target_mem_desc *gomp_map_vars_openacc (struct gomp_device_descr *,
+						      struct goacc_asyncqueue *,
+						      size_t, void **, size_t *,
+						      unsigned short *, void *);
 extern void gomp_unmap_tgt (struct target_mem_desc *);
 extern void gomp_unmap_vars (struct target_mem_desc *, bool);
 extern void gomp_unmap_vars_async (struct target_mem_desc *, bool,
Index: libgomp/libgomp_g.h
===================================================================
--- libgomp/libgomp_g.h	(revision 278656)
+++ libgomp/libgomp_g.h	(working copy)
@@ -364,7 +364,7 @@  extern void GOACC_parallel_keyed (int, void (*) (v
 extern void GOACC_parallel (int, void (*) (void *), size_t, void **, size_t *,
 			    unsigned short *, int, int, int, int, int, ...);
 extern void GOACC_data_start (int, size_t, void **, size_t *,
-			      unsigned short *);
+			      unsigned short *, ...);
 extern void GOACC_data_end (void);
 extern void GOACC_enter_exit_data (int, size_t, void **,
 				   size_t *, unsigned short *, int, int, ...);
Index: libgomp/oacc-int.h
===================================================================
--- libgomp/oacc-int.h	(revision 278656)
+++ libgomp/oacc-int.h	(working copy)
@@ -164,6 +164,57 @@  bool _goacc_profiling_setup_p (struct goacc_thread
 void goacc_profiling_dispatch (acc_prof_info *, acc_event_info *,
 			       acc_api_info *);
 
+/* Definitions for data structures describing OpenACC non-contiguous arrays
+   (Note: interfaces with compiler)
+
+   The compiler generates a descriptor for each such array, places the
+   descriptor on stack, and passes the address of the descriptor to the libgomp
+   runtime as a normal map argument. The runtime then processes the array
+   data structure setup, and replaces the argument with the new actual
+   array address for the child function.
+
+   Care must be taken such that the struct field and layout assumptions
+   of struct goacc_ncarray_dim, goacc_ncarray_descr_type inside the compiler
+   be consistant with the below declarations.  */
+
+struct goacc_ncarray_dim {
+  size_t base;
+  size_t length;
+  size_t elem_size;
+  size_t is_array;
+};
+
+struct goacc_ncarray_descr_type
+{
+  size_t ndims;
+  struct goacc_ncarray_dim dims[];
+};
+
+/* Internal non-contiguous array info struct, used only here inside the runtime. */
+
+struct goacc_ncarray
+{
+  struct goacc_ncarray_descr_type *descr;
+  void *ptr;
+  size_t map_index;
+  size_t ptrblock_size;
+  void **data_rows;
+  void **tgt_data_rows;
+  size_t data_row_num;
+  size_t data_row_size;
+};
+
+struct goacc_ncarray_info
+{
+  size_t num_data_rows, num_ncarray;
+  void **data_rows;
+  void **tgt_data_rows;
+  struct goacc_ncarray ncarray[];
+};
+
+extern void *goacc_noncontig_array_create_ptrblock (struct goacc_ncarray *, void *);
+
+
 #ifdef HAVE_ATTRIBUTE_VISIBILITY
 # pragma GCC visibility pop
 #endif
Index: libgomp/oacc-parallel.c
===================================================================
--- libgomp/oacc-parallel.c	(revision 278656)
+++ libgomp/oacc-parallel.c	(working copy)
@@ -36,8 +36,8 @@ 
 #include <string.h>
 #include <stdarg.h>
 #include <assert.h>
+#include <stdio.h>
 
-
 /* In the ABI, the GOACC_FLAGs are encoded as an inverted bitmask, so that we
    continue to support the following two legacy values.  */
 _Static_assert (GOACC_FLAGS_UNMARSHAL (GOMP_DEVICE_ICV) == 0,
@@ -113,7 +113,174 @@  handle_ftn_pointers (size_t mapnum, void **hostadd
 
 static void goacc_wait (int async, int num_waits, va_list *ap);
 
+static size_t
+goacc_noncontig_array_count_rows (struct goacc_ncarray_descr_type *descr)
+{
+  size_t nrows = 1;
+  for (size_t d = 0; d < descr->ndims - 1; d++)
+    nrows *= descr->dims[d].length / sizeof (void *);
+  return nrows;
+}
 
+static void
+goacc_noncontig_array_compute_sizes (struct goacc_ncarray *nca)
+{
+  size_t d, n = 1;
+  struct goacc_ncarray_descr_type *descr = nca->descr;
+
+  nca->ptrblock_size = 0;
+  for (d = 0; d < descr->ndims - 1; d++)
+    {
+      size_t dim_count = descr->dims[d].length / descr->dims[d].elem_size;
+      size_t dim_ptrblock_size = (descr->dims[d + 1].is_array
+				  ? 0 : descr->dims[d].length * n);
+      nca->ptrblock_size += dim_ptrblock_size;
+      n *= dim_count;
+    }
+  nca->data_row_num = n;
+  nca->data_row_size = descr->dims[d].length;
+  
+}
+
+static void
+goacc_noncontig_array_fill_rows_1 (struct goacc_ncarray_descr_type *descr, void *nca,
+				   size_t d, void ***row_ptr, size_t *count)
+{
+  if (d < descr->ndims - 1)
+    {
+      size_t elsize = descr->dims[d].elem_size;
+      size_t n = descr->dims[d].length / elsize;
+      void *p = nca + descr->dims[d].base;
+      for (size_t i = 0; i < n; i++)
+	{
+	  void *ptr = p + i * elsize;
+	  /* Deref if next dimension is not array.  */
+	  if (!descr->dims[d + 1].is_array)
+	    ptr = *((void **) ptr);
+	  goacc_noncontig_array_fill_rows_1 (descr, ptr, d + 1, row_ptr, count);
+	}
+    }
+  else
+    {
+      **row_ptr = nca + descr->dims[d].base;
+      *row_ptr += 1;
+      *count += 1;
+    }
+}
+
+static size_t
+goacc_noncontig_array_fill_rows (struct goacc_ncarray *nca)
+{
+  size_t count = 0;
+  void **p = nca->data_rows;
+  goacc_noncontig_array_fill_rows_1 (nca->descr, nca->ptr, 0, &p, &count);
+  return count;
+}
+
+static struct goacc_ncarray_info *
+goacc_process_noncontiguous_arrays (size_t mapnum, void **hostaddrs,
+				    unsigned short *kinds, va_list* ap)
+{
+  size_t i, nr, num_data_rows = 0, num_ncarray = 0, curr_row_start = 0;
+  struct goacc_ncarray_descr_type *descr;
+
+  /* We need to go over *ap twice, so preserve *ap state here.  */
+  va_list itr;
+  va_copy (itr, *ap);
+  for (i = 0; i < mapnum; i++)
+    if (GOMP_MAP_NONCONTIG_ARRAY_P (kinds[i] & 0xff))
+      {
+	descr = va_arg (itr, struct goacc_ncarray_descr_type *);
+	num_data_rows += goacc_noncontig_array_count_rows (descr);
+	num_ncarray += 1;
+      }
+    else
+      break;
+
+  /* Allocate the entire info struct, array entries, and row pointer
+     arrays in one large block.  */
+  struct goacc_ncarray_info *nca_info
+    = gomp_malloc (sizeof (struct goacc_ncarray_info)
+		   + sizeof (struct goacc_ncarray) * num_ncarray
+		   + sizeof (void *) * num_data_rows * 2);
+  nca_info->num_data_rows = num_data_rows;
+  nca_info->num_ncarray = num_ncarray;
+  nca_info->data_rows = (void **) (nca_info->ncarray + num_ncarray);
+  nca_info->tgt_data_rows = nca_info->data_rows + num_data_rows;
+
+  struct goacc_ncarray *curr_ncarray = nca_info->ncarray;
+  for (i = 0; i < mapnum; i++)
+    if (GOMP_MAP_NONCONTIG_ARRAY_P (kinds[i] & 0xff))
+      {
+	descr = va_arg (*ap, struct goacc_ncarray_descr_type *);
+	curr_ncarray->descr = descr;
+	curr_ncarray->ptr = hostaddrs[i];
+	curr_ncarray->map_index = i;
+
+	goacc_noncontig_array_compute_sizes (curr_ncarray);
+
+	curr_ncarray->data_rows = nca_info->data_rows + curr_row_start;
+	curr_ncarray->tgt_data_rows = nca_info->tgt_data_rows + curr_row_start;
+
+	nr = goacc_noncontig_array_fill_rows (curr_ncarray);
+	assert (nr == curr_ncarray->data_row_num);
+	curr_row_start += nr;
+	curr_ncarray += 1;
+      }
+    else
+      break;
+
+  return nca_info;
+}
+
+void *
+goacc_noncontig_array_create_ptrblock (struct goacc_ncarray *nca,
+				       void *tgt_ptrblock_addr)
+{
+  struct goacc_ncarray_descr_type *descr = nca->descr;
+  void **tgt_data_rows = nca->tgt_data_rows;
+  void *ptrblock = gomp_malloc (nca->ptrblock_size);
+  void **curr_dim_ptrblock = (void **) ptrblock;
+  size_t n = 1;
+
+  for (size_t d = 0; d < descr->ndims - 1; d++)
+    {
+      int curr_dim_len = descr->dims[d].length;
+      int next_dim_len = descr->dims[d + 1].length;
+      int curr_dim_num = curr_dim_len / sizeof (void *);
+      size_t next_dim_bias = descr->dims[d + 1].base;
+
+      void *next_dim_ptrblock
+	= (void *)(curr_dim_ptrblock + n * curr_dim_num);
+
+      for (int b = 0; b < n; b++)
+        for (int i = 0; i < curr_dim_num; i++)
+	  {
+	    if (d < descr->ndims - 2)
+	      {
+		void *ptr = (next_dim_ptrblock
+			     + b * curr_dim_num * next_dim_len
+			     + i * next_dim_len);
+		void *tgt_ptr = (tgt_ptrblock_addr
+				 + (ptr - ptrblock) - next_dim_bias);
+		curr_dim_ptrblock[b * curr_dim_num + i] = tgt_ptr;
+	      }
+	    else
+	      {
+		curr_dim_ptrblock[b * curr_dim_num + i]
+		  = tgt_data_rows[b * curr_dim_num + i] - next_dim_bias;
+	      }
+	    void *addr = &curr_dim_ptrblock[b * curr_dim_num + i];
+	    assert (ptrblock <= addr && addr < ptrblock + nca->ptrblock_size);
+	  }
+
+      n *= curr_dim_num;
+      curr_dim_ptrblock = next_dim_ptrblock;
+    }
+  assert (n == nca->data_row_num);
+  return ptrblock;
+}
+
 /* Launch a possibly offloaded function with FLAGS.  FN is the host fn
    address.  MAPNUM, HOSTADDRS, SIZES & KINDS  describe the memory
    blocks to be copied to/from the device.  Varadic arguments are
@@ -138,6 +305,7 @@  GOACC_parallel_keyed (int flags_m, void (*fn) (voi
   int async = GOMP_ASYNC_SYNC;
   unsigned dims[GOMP_DIM_MAX];
   unsigned tag;
+  struct goacc_ncarray_info *nca_info = NULL;
 
 #ifdef HAVE_INTTYPES_H
   gomp_debug (0, "%s: mapnum=%"PRIu64", hostaddrs=%p, size=%p, kinds=%p\n",
@@ -270,11 +438,20 @@  GOACC_parallel_keyed (int flags_m, void (*fn) (voi
 	    break;
 	  }
 
+	  /*case GOMP_LAUNCH_NONCONTIG_ARRAYS:
+	  nca_info = goacc_process_noncontiguous_arrays (mapnum, hostaddrs,
+							 kinds, &ap);
+							 break;*/
+
 	default:
 	  gomp_fatal ("unrecognized offload code '%d',"
 		      " libgomp is too old", GOMP_LAUNCH_CODE (tag));
 	}
     }
+
+  if (mapnum > 0 && GOMP_MAP_NONCONTIG_ARRAY_P (kinds[0] & 0xff))
+    nca_info = goacc_process_noncontiguous_arrays (mapnum, hostaddrs, kinds, &ap);
+  
   va_end (ap);
   
   if (!(acc_dev->capabilities & GOMP_OFFLOAD_CAP_NATIVE_EXEC))
@@ -311,8 +488,10 @@  GOACC_parallel_keyed (int flags_m, void (*fn) (voi
 
   goacc_aq aq = get_goacc_asyncqueue (async);
 
-  tgt = gomp_map_vars_async (acc_dev, aq, mapnum, hostaddrs, NULL, sizes, kinds,
-			     true, GOMP_MAP_VARS_OPENACC);
+  tgt = gomp_map_vars_openacc (acc_dev, aq, mapnum, hostaddrs, sizes, kinds,
+			       nca_info);
+  free (nca_info);
+
   if (profiling_p)
     {
       prof_info.event_type = acc_ev_enter_data_end;
@@ -390,7 +569,7 @@  GOACC_parallel (int flags_m, void (*fn) (void *),
 
 void
 GOACC_data_start (int flags_m, size_t mapnum,
-		  void **hostaddrs, size_t *sizes, unsigned short *kinds)
+		  void **hostaddrs, size_t *sizes, unsigned short *kinds, ...)
 {
   int flags = GOACC_FLAGS_UNMARSHAL (flags_m);
 
@@ -480,8 +659,7 @@  GOACC_data_start (int flags_m, size_t mapnum,
     {
       prof_info.device_type = acc_device_host;
       api_info.device_type = prof_info.device_type;
-      tgt = gomp_map_vars (NULL, 0, NULL, NULL, NULL, NULL, true,
-			   GOMP_MAP_VARS_OPENACC);
+      tgt = gomp_map_vars_openacc (NULL, NULL, 0, NULL, NULL, NULL, NULL);
       tgt->prev = thr->mapped_data;
       thr->mapped_data = tgt;
 
@@ -488,9 +666,19 @@  GOACC_data_start (int flags_m, size_t mapnum,
       goto out_prof;
     }
 
+  struct goacc_ncarray_info *nca_info = NULL;
+  if (mapnum > 0 && GOMP_MAP_NONCONTIG_ARRAY_P (kinds[0] & 0xff))
+    {
+      va_list ap;
+      va_start (ap, kinds);
+      nca_info = goacc_process_noncontiguous_arrays (mapnum, hostaddrs, kinds, &ap);
+      va_end (ap);
+    }
+  
   gomp_debug (0, "  %s: prepare mappings\n", __FUNCTION__);
-  tgt = gomp_map_vars (acc_dev, mapnum, hostaddrs, NULL, sizes, kinds, true,
-		       GOMP_MAP_VARS_OPENACC);
+  tgt = gomp_map_vars_openacc (acc_dev, NULL, mapnum, hostaddrs, sizes, kinds,
+			       nca_info);
+  free (nca_info);
   gomp_debug (0, "  %s: mappings prepared\n", __FUNCTION__);
   tgt->prev = thr->mapped_data;
   thr->mapped_data = tgt;
@@ -701,7 +889,6 @@  GOACC_enter_exit_data (int flags_m, size_t mapnum,
   if (num_waits)
     {
       va_list ap;
-
       va_start (ap, num_waits);
       goacc_wait (async, num_waits, &ap);
       va_end (ap);
Index: libgomp/target.c
===================================================================
--- libgomp/target.c	(revision 278656)
+++ libgomp/target.c	(working copy)
@@ -524,10 +524,11 @@  static inline __attribute__((always_inline)) struc
 gomp_map_vars_internal (struct gomp_device_descr *devicep,
 			struct goacc_asyncqueue *aq, size_t mapnum,
 			void **hostaddrs, void **devaddrs, size_t *sizes,
-			void *kinds, bool short_mapkind,
-			enum gomp_map_vars_kind pragma_kind)
+			void *kinds, struct goacc_ncarray_info *nca_info,
+			bool short_mapkind, enum gomp_map_vars_kind pragma_kind)
 {
   size_t i, tgt_align, tgt_size, not_found_cnt = 0;
+  size_t nca_data_row_num = (nca_info ? nca_info->num_data_rows : 0);
   bool has_firstprivate = false;
   const int rshift = short_mapkind ? 8 : 3;
   const int typemask = short_mapkind ? 0xff : 0x7;
@@ -534,8 +535,9 @@  gomp_map_vars_internal (struct gomp_device_descr *
   struct splay_tree_s *mem_map = &devicep->mem_map;
   struct splay_tree_key_s cur_node;
   struct target_mem_desc *tgt
-    = gomp_malloc (sizeof (*tgt) + sizeof (tgt->list[0]) * mapnum);
-  tgt->list_count = mapnum;
+    = gomp_malloc (sizeof (*tgt)
+		   + sizeof (tgt->list[0]) * (mapnum + nca_data_row_num));
+  tgt->list_count = mapnum + nca_data_row_num;
   tgt->refcount = pragma_kind == GOMP_MAP_VARS_ENTER_DATA ? 0 : 1;
   tgt->device_descr = devicep;
   struct gomp_coalesce_buf cbuf, *cbufp = NULL;
@@ -667,6 +669,27 @@  gomp_map_vars_internal (struct gomp_device_descr *
 	  has_firstprivate = true;
 	  continue;
 	}
+      else if (GOMP_MAP_NONCONTIG_ARRAY_P (kind & typemask))
+	{
+	  /* Ignore non-contiguous arrays for now, we process them together
+	     later.  */
+	  tgt->list[i].key = NULL;
+	  tgt->list[i].offset = 0;
+	  not_found_cnt++;
+
+	  /* The map for the non-contiguous array itself is never copied from
+	     during unmapping, its the data rows that count. Set copy-from
+	     flags to false here.  */
+	  tgt->list[i].copy_from = false;
+	  tgt->list[i].always_copy_from = false;
+
+	  size_t align = (size_t) 1 << (kind >> rshift);
+	  if (tgt_align < align)
+	    tgt_align = align;
+
+	  continue;
+	}
+      
       cur_node.host_start = (uintptr_t) hostaddrs[i];
       if (!GOMP_MAP_POINTER_P (kind & typemask))
 	cur_node.host_end = cur_node.host_start + sizes[i];
@@ -735,6 +758,44 @@  gomp_map_vars_internal (struct gomp_device_descr *
 	}
     }
 
+  /* For non-contiguous arrays. Each data row is one target item, separated
+     from the normal map clause items, hence we order them after mapnum.  */
+  if (nca_info)
+    {
+      struct target_var_desc *next_var_desc = &tgt->list[mapnum];
+      for (i = 0; i < nca_info->num_ncarray; i++)
+	{
+	  struct goacc_ncarray *nca = &nca_info->ncarray[i];
+	  int kind = get_kind (short_mapkind, kinds, nca->map_index);
+	  size_t align = (size_t) 1 << (kind >> rshift);
+	  tgt_size = (tgt_size + align - 1) & ~(align - 1);
+	  tgt_size += nca->ptrblock_size;
+	  
+	  for (size_t j = 0; j < nca->data_row_num; j++)
+	    {
+	      struct target_var_desc *row_desc = next_var_desc++;
+	      void *row = nca->data_rows[j];
+	      cur_node.host_start = (uintptr_t) row;
+	      cur_node.host_end = cur_node.host_start + nca->data_row_size;
+	      splay_tree_key n = splay_tree_lookup (mem_map, &cur_node);
+	      if (n)
+		{
+		  assert (n->refcount != REFCOUNT_LINK);
+		  gomp_map_vars_existing (devicep, aq, n, &cur_node, row_desc,
+					  kind & typemask,
+					  /* TODO: cbuf? */ NULL);
+		}
+	      else
+		{
+		  tgt_size = (tgt_size + align - 1) & ~(align - 1);
+		  tgt_size += nca->data_row_size;
+		  not_found_cnt++;
+		}
+	    }
+	}
+      assert (next_var_desc == &tgt->list[mapnum + nca_info->num_data_rows]);
+    }
+
   if (devaddrs)
     {
       if (mapnum != 1)
@@ -895,6 +956,15 @@  gomp_map_vars_internal (struct gomp_device_descr *
 	      default:
 		break;
 	      }
+
+	    if (GOMP_MAP_NONCONTIG_ARRAY_P (kind & typemask))
+	      {
+		tgt->list[i].key = &array->key;
+		tgt->list[i].key->tgt = tgt;
+		array++;
+		continue;
+	      }
+
 	    splay_tree_key k = &array->key;
 	    k->host_start = (uintptr_t) hostaddrs[i];
 	    if (!GOMP_MAP_POINTER_P (kind & typemask))
@@ -1044,6 +1114,98 @@  gomp_map_vars_internal (struct gomp_device_descr *
 		array++;
 	      }
 	  }
+
+      /* Processing of non-contiguous array rows.  */
+      if (nca_info)
+	{
+	  struct target_var_desc *next_var_desc = &tgt->list[mapnum];
+	  for (i = 0; i < nca_info->num_ncarray; i++)
+	    {
+	      struct goacc_ncarray *nca = &nca_info->ncarray[i];
+	      int kind = get_kind (short_mapkind, kinds, nca->map_index);
+	      size_t align = (size_t) 1 << (kind >> rshift);
+	      tgt_size = (tgt_size + align - 1) & ~(align - 1);
+
+	      assert (nca->ptr == hostaddrs[nca->map_index]);
+
+	      /* For the map of the non-contiguous array itself, adjust so that
+		 the passed device address points to the beginning of the
+		 ptrblock. Remember to adjust the first-dimension's bias here.   */
+	      tgt->list[nca->map_index].key->tgt_offset
+		= tgt_size - nca->descr->dims[0].base;
+
+	      void *target_ptrblock = (void*) tgt->tgt_start + tgt_size;
+	      tgt_size += nca->ptrblock_size;
+
+	      /* Add splay key for each data row in current non-contiguous
+		 array.  */
+	      for (size_t j = 0; j < nca->data_row_num; j++)
+		{
+		  struct target_var_desc *row_desc = next_var_desc++;
+		  void *row = nca->data_rows[j];
+		  cur_node.host_start = (uintptr_t) row;
+		  cur_node.host_end = cur_node.host_start + nca->data_row_size;
+		  splay_tree_key k = splay_tree_lookup (mem_map, &cur_node);
+		  if (k)
+		    {
+		      assert (k->refcount != REFCOUNT_LINK);
+		      gomp_map_vars_existing (devicep, aq, k, &cur_node, row_desc,
+					      kind & typemask,
+					      cbufp);
+		    }
+		  else
+		    {
+		      tgt->refcount++;
+		      tgt_size = (tgt_size + align - 1) & ~(align - 1);
+
+		      k = &array->key;
+		      k->host_start = (uintptr_t) row;
+		      k->host_end = k->host_start + nca->data_row_size;
+
+		      k->tgt = tgt;
+		      k->refcount = 1;
+		      k->dynamic_refcount = 0;
+		      k->link_key = NULL;
+		      k->tgt_offset = tgt_size;
+
+		      tgt_size += nca->data_row_size;
+		      
+		      row_desc->key = k;
+		      row_desc->copy_from
+			= GOMP_MAP_COPY_FROM_P (kind & typemask);
+		      row_desc->always_copy_from
+			= GOMP_MAP_COPY_FROM_P (kind & typemask);
+		      row_desc->offset = 0;
+		      row_desc->length = nca->data_row_size;
+
+		      array->left = NULL;
+		      array->right = NULL;
+		      splay_tree_insert (mem_map, array);
+
+		      if (GOMP_MAP_COPY_TO_P (kind & typemask))
+			gomp_copy_host2dev (devicep, aq,
+					    (void *) tgt->tgt_start + k->tgt_offset,
+					    (void *) k->host_start,
+					    nca->data_row_size, cbufp);
+		      array++;
+		    }
+		  nca->tgt_data_rows[j]
+		    = (void *) (k->tgt->tgt_start + k->tgt_offset);
+		}
+
+	      /* Now we have the target memory allocated, and target offsets of all
+		 row blocks assigned and calculated, we can construct the
+		 accelerator side ptrblock and copy it in.  */
+	      if (nca->ptrblock_size)
+		{
+		  void *ptrblock = goacc_noncontig_array_create_ptrblock
+		    (nca, target_ptrblock);
+		  gomp_copy_host2dev (devicep, aq, target_ptrblock, ptrblock,
+				      nca->ptrblock_size, cbufp);
+		  free (ptrblock);
+		}
+	    }
+	}
     }
 
   if (pragma_kind == GOMP_MAP_VARS_TARGET)
@@ -1086,12 +1248,25 @@  gomp_map_vars_internal (struct gomp_device_descr *
 }
 
 attribute_hidden struct target_mem_desc *
+gomp_map_vars_openacc (struct gomp_device_descr *devicep,
+		       struct goacc_asyncqueue *aq, size_t mapnum,
+		       void **hostaddrs, size_t *sizes, unsigned short *kinds,
+		       void *nca_info)
+{
+  return gomp_map_vars_internal (devicep, aq, mapnum, hostaddrs, NULL,
+				 sizes, (void *) kinds,
+				 (struct goacc_ncarray_info *) nca_info,
+				 true, GOMP_MAP_VARS_OPENACC);
+}
+
+attribute_hidden struct target_mem_desc *
 gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
 	       void **hostaddrs, void **devaddrs, size_t *sizes, void *kinds,
 	       bool short_mapkind, enum gomp_map_vars_kind pragma_kind)
 {
   return gomp_map_vars_internal (devicep, NULL, mapnum, hostaddrs, devaddrs,
-				 sizes, kinds, short_mapkind, pragma_kind);
+				 sizes, kinds, NULL, short_mapkind,
+				 pragma_kind);
 }
 
 attribute_hidden struct target_mem_desc *
@@ -1102,7 +1277,7 @@  gomp_map_vars_async (struct gomp_device_descr *dev
 		     enum gomp_map_vars_kind pragma_kind)
 {
   return gomp_map_vars_internal (devicep, aq, mapnum, hostaddrs, devaddrs,
-				 sizes, kinds, short_mapkind, pragma_kind);
+				 sizes, kinds, NULL, short_mapkind, pragma_kind);
 }
 
 attribute_hidden void
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-1.c
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-1.c	(nonexistent)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-1.c	(working copy)
@@ -0,0 +1,103 @@ 
+/* { dg-do run } */
+
+#include <stdlib.h>
+#include <assert.h>
+
+#define n 100
+#define m 100
+
+int b[n][m];
+
+void
+test1 (void)
+{
+  int i, j, *a[100];
+
+  /* Array of pointers form test.  */
+  for (i = 0; i < n; i++)
+    {
+      a[i] = (int *)malloc (sizeof (int) * m);
+      for (j = 0; j < m; j++)
+	b[i][j] = j - i;
+    }
+
+  #pragma acc parallel loop copyout(a[0:n][0:m]) copyin(b)
+  for (i = 0; i < n; i++)
+    #pragma acc loop
+    for (j = 0; j < m; j++)
+      a[i][j] = b[i][j];
+
+  for (i = 0; i < n; i++)
+    {
+      for (j = 0; j < m; j++)
+	assert (a[i][j] == b[i][j]);
+      /* Clean up.  */
+      free (a[i]);
+    }
+}
+
+void
+test2 (void)
+{
+  int i, j, **a = (int **) malloc (sizeof (int *) * n);
+
+  /* Separately allocated blocks.  */
+  for (i = 0; i < n; i++)
+    {
+      a[i] = (int *)malloc (sizeof (int) * m);
+      for (j = 0; j < m; j++)
+	b[i][j] = j - i;
+    }
+
+  #pragma acc parallel loop copyout(a[0:n][0:m]) copyin(b)
+  for (i = 0; i < n; i++)
+    #pragma acc loop
+    for (j = 0; j < m; j++)
+      a[i][j] = b[i][j];
+
+  for (i = 0; i < n; i++)
+    {
+      for (j = 0; j < m; j++)
+	assert (a[i][j] == b[i][j]);
+      /* Clean up.  */
+      free (a[i]);
+    }
+  free (a);
+}
+
+void
+test3 (void)
+{
+  int i, j, **a = (int **) malloc (sizeof (int *) * n);
+  a[0] = (int *) malloc (sizeof (int) * n * m);
+
+  /* Rows allocated in one contiguous block.  */
+  for (i = 0; i < n; i++)
+    {
+      a[i] = *a + i * m;
+      for (j = 0; j < m; j++)
+	b[i][j] = j - i;
+    }
+
+  #pragma acc parallel loop copyout(a[0:n][0:m]) copyin(b)
+  for (i = 0; i < n; i++)
+    #pragma acc loop
+    for (j = 0; j < m; j++)
+      a[i][j] = b[i][j];
+
+  for (i = 0; i < n; i++)
+    for (j = 0; j < m; j++)
+      assert (a[i][j] == b[i][j]);
+
+  free (a[0]);
+  free (a);
+}
+
+int
+main (void)
+{
+  test1 ();
+  test2 ();
+  test3 ();
+  return 0;
+}
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-2.c
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-2.c	(nonexistent)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-2.c	(working copy)
@@ -0,0 +1,37 @@ 
+/* { dg-do run } */
+
+#include <assert.h>
+#include "noncontig_array-utils.h"
+
+int
+main (void)
+{
+  int n = 10;
+  int ***a = (int ***) create_ncarray (sizeof (int), n, 3);
+  int ***b = (int ***) create_ncarray (sizeof (int), n, 3);
+  int ***c = (int ***) create_ncarray (sizeof (int), n, 3);
+
+  for (int i = 0; i < n; i++)
+    for (int j = 0; j < n; j++)
+      for (int k = 0; k < n; k++)
+	{
+	  a[i][j][k] = i + j * k + k;
+	  b[i][j][k] = j + k * i + i * j;
+	  c[i][j][k] = a[i][j][k];
+	}
+
+  #pragma acc parallel copy (a[0:n][0:n][0:n]) copyin (b[0:n][0:n][0:n])
+  {
+    for (int i = 0; i < n; i++)
+      for (int j = 0; j < n; j++)
+	for (int k = 0; k < n; k++)
+	  a[i][j][k] += b[k][j][i] + i + j + k;
+  }
+
+  for (int i = 0; i < n; i++)
+    for (int j = 0; j < n; j++)
+      for (int k = 0; k < n; k++)
+	assert (a[i][j][k] == c[i][j][k] + b[k][j][i] + i + j + k);
+
+  return 0;
+}
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-3.c
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-3.c	(nonexistent)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-3.c	(working copy)
@@ -0,0 +1,45 @@ 
+/* { dg-do run } */
+
+#include <assert.h>
+#include "noncontig_array-utils.h"
+
+int main (void)
+{
+  int n = 20, x = 5, y = 12;
+  int *****a = (int *****) create_ncarray (sizeof (int), n, 5);
+
+  int sum1 = 0, sum2 = 0, sum3 = 0;
+
+  for (int i = 0; i < n; i++)
+    for (int j = 0; j < n; j++)
+      for (int k = 0; k < n; k++)
+	for (int l = 0; l < n; l++)
+	  for (int m = 0; m < n; m++)
+	    {
+	      a[i][j][k][l][m] = 1;
+	      sum1++;
+	    }
+
+  #pragma acc parallel copy (a[x:y][x:y][x:y][x:y][x:y]) copy(sum2)
+  {
+    for (int i = x; i < x + y; i++)
+      for (int j = x; j < x + y; j++)
+	for (int k = x; k < x + y; k++)
+	  for (int l = x; l < x + y; l++)
+	    for (int m = x; m < x + y; m++)
+	      {
+		a[i][j][k][l][m] = 0;
+		sum2++;
+	      }
+  }
+
+  for (int i = 0; i < n; i++)
+    for (int j = 0; j < n; j++)
+      for (int k = 0; k < n; k++)
+	for (int l = 0; l < n; l++)
+	  for (int m = 0; m < n; m++)
+	    sum3 += a[i][j][k][l][m];
+
+  assert (sum1 == sum2 + sum3);
+  return 0;
+}
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-4.c
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-4.c	(nonexistent)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-4.c	(working copy)
@@ -0,0 +1,36 @@ 
+/* { dg-do run } */
+
+#include <assert.h>
+#include "noncontig_array-utils.h"
+
+int main (void)
+{
+  int n = 128;
+  double ***a = (double ***) create_ncarray (sizeof (double), n, 3);
+  double ***b = (double ***) create_ncarray (sizeof (double), n, 3);
+
+  for (int i = 0; i < n; i++)
+    for (int j = 0; j < n; j++)
+      for (int k = 0; k < n; k++)
+	a[i][j][k] = i + j + k + i * j * k;
+
+  /* This test exercises async copyout of non-contiguous array rows.  */
+  #pragma acc parallel copyin(a[0:n][0:n][0:n]) copyout(b[0:n][0:n][0:n]) async(5)
+  {
+    #pragma acc loop gang
+    for (int i = 0; i < n; i++)
+      #pragma acc loop vector
+      for (int j = 0; j < n; j++)
+	for (int k = 0; k < n; k++)
+	  b[i][j][k] = a[i][j][k] * 2.0;
+  }
+
+  #pragma acc wait (5)
+
+  for (int i = 0; i < n; i++)
+    for (int j = 0; j < n; j++)
+      for (int k = 0; k < n; k++)
+	assert (b[i][j][k] == a[i][j][k] * 2.0);
+
+  return 0;
+}
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-utils.h
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-utils.h	(nonexistent)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-utils.h	(working copy)
@@ -0,0 +1,44 @@ 
+#include <stdlib.h>
+#include <string.h>
+#include <assert.h>
+#include <stdint.h>
+
+/* Allocate and create a pointer based NDIMS-dimensional array,
+   each dimension DIMLEN long, with ELSIZE sized data elements.  */
+void *
+create_ncarray (size_t elsize, int dimlen, int ndims)
+{
+  size_t blk_size = 0;
+  size_t n = 1;
+
+  for (int i = 0; i < ndims - 1; i++)
+    {
+      n *= dimlen;
+      blk_size += sizeof (void *) * n;
+    }
+  size_t data_rows_num = n;
+  size_t data_rows_offset = blk_size;
+  blk_size += elsize * n * dimlen;
+
+  void *blk = (void *) malloc (blk_size);
+  memset (blk, 0, blk_size);
+  void **curr_dim = (void **) blk;
+  n = 1;
+
+  for (int d = 0; d < ndims - 1; d++)
+    {
+      uintptr_t next_dim = (uintptr_t) (curr_dim + n * dimlen);
+      size_t next_dimlen = dimlen * (d < ndims - 2 ? sizeof (void *) : elsize);
+
+      for (int b = 0; b < n; b++)
+        for (int i = 0; i < dimlen; i++)
+	  if (d < ndims - 1)
+	    curr_dim[b * dimlen + i]
+	      = (void*) (next_dim + b * dimlen * next_dimlen + i * next_dimlen);
+
+      n *= dimlen;
+      curr_dim = (void**) next_dim;
+    }
+  assert (n == data_rows_num);
+  return blk;
+}