diff --git a/gcc/c-family/c-common.h b/gcc/c-family/c-common.h
index ded1b5ded04..9177a9342db 100644
--- a/gcc/c-family/c-common.h
+++ b/gcc/c-family/c-common.h
@@ -1275,8 +1275,10 @@ enum c_omp_region_type
   C_ORT_ACC			= 1 << 1,
   C_ORT_DECLARE_SIMD		= 1 << 2,
   C_ORT_TARGET			= 1 << 3,
+  C_ORT_EXIT_DATA		= 1 << 4,
   C_ORT_OMP_DECLARE_SIMD	= C_ORT_OMP | C_ORT_DECLARE_SIMD,
   C_ORT_OMP_TARGET		= C_ORT_OMP | C_ORT_TARGET,
+  C_ORT_OMP_EXIT_DATA		= C_ORT_OMP | C_ORT_EXIT_DATA,
   C_ORT_ACC_TARGET		= C_ORT_ACC | C_ORT_TARGET
 };
 
@@ -1362,11 +1364,6 @@ public:
   virtual bool check_clause (tree);
   tree get_root_term (bool);
 
-  tree get_address ()
-    {
-      return orig;
-    }
-
   tree unconverted_ref_origin ();
   bool component_access_p ();
 
@@ -1378,9 +1375,9 @@ public:
   bool maybe_zero_length_array_section (tree);
 
   tree expand_array_base (tree, vec<omp_addr_token *> &, tree, unsigned *,
-			  c_omp_region_type, bool);
+			  c_omp_region_type);
   tree expand_component_selector (tree, vec<omp_addr_token *> &, tree,
-				  unsigned *);
+				  unsigned *, c_omp_region_type);
   tree expand_map_clause (tree, tree, vec<omp_addr_token *> &,
 			  c_omp_region_type);
 };
diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc
index 6700567afaf..5e534aa21cd 100644
--- a/gcc/c-family/c-omp.cc
+++ b/gcc/c-family/c-omp.cc
@@ -3447,6 +3447,13 @@ c_omp_address_inspector::maybe_zero_length_array_section (tree clause)
     case GOMP_MAP_ALWAYS_TO:
     case GOMP_MAP_ALWAYS_FROM:
     case GOMP_MAP_ALWAYS_TOFROM:
+    case GOMP_MAP_PRESENT_ALLOC:
+    case GOMP_MAP_PRESENT_TO:
+    case GOMP_MAP_PRESENT_FROM:
+    case GOMP_MAP_PRESENT_TOFROM:
+    case GOMP_MAP_ALWAYS_PRESENT_TO:
+    case GOMP_MAP_ALWAYS_PRESENT_FROM:
+    case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
     case GOMP_MAP_RELEASE:
     case GOMP_MAP_DELETE:
     case GOMP_MAP_FORCE_TO:
@@ -3461,26 +3468,31 @@ c_omp_address_inspector::maybe_zero_length_array_section (tree clause)
 
 /* Expand a chained access.  We only expect to see a quite limited range of
    expression types here, because e.g. you can't have an array of
-   references.  See also gimplify.cc:omp_expand_access_chain.  */
+   references.  */
 
 static tree
 omp_expand_access_chain (tree c, tree expr, vec<omp_addr_token *> &addr_tokens,
-			 unsigned *idx)
+			 unsigned *idx, c_omp_region_type ort)
 {
   using namespace omp_addr_tokenizer;
   location_t loc = OMP_CLAUSE_LOCATION (c);
   unsigned i = *idx;
   tree c2 = NULL_TREE;
-  gomp_map_kind kind
-    = (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM
-       || (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
-	   && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FROM
-	       || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DELETE
-	       || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_RELEASE
-	       || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_FROM
-	       || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_FROM
-	       || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PRESENT_FROM)))
-      ? GOMP_MAP_DETACH : GOMP_MAP_ATTACH;
+  gomp_map_kind kind;
+
+  if ((ort & C_ORT_EXIT_DATA) != 0
+      || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM
+      || (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+	  && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FROM
+	      || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DELETE
+	      || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_RELEASE
+	      || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_FROM
+	      || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_FROM
+	      || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PRESENT_FROM
+	      || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_PRESENT_FROM)))
+    kind = GOMP_MAP_DETACH;
+  else
+    kind = GOMP_MAP_ATTACH;
 
   switch (addr_tokens[i]->u.access_kind)
     {
@@ -3519,7 +3531,7 @@ omp_expand_access_chain (tree c, tree expr, vec<omp_addr_token *> &addr_tokens,
 
   if (i < addr_tokens.length ()
       && addr_tokens[i]->type == ACCESS_METHOD)
-    return omp_expand_access_chain (c, expr, addr_tokens, idx);
+    return omp_expand_access_chain (c, expr, addr_tokens, idx, ort);
 
   return c;
 }
@@ -3530,41 +3542,41 @@ tree
 c_omp_address_inspector::expand_array_base (tree c,
 					    vec<omp_addr_token *> &addr_tokens,
 					    tree expr, unsigned *idx,
-					    c_omp_region_type ort,
-					    bool decl_p)
+					    c_omp_region_type ort)
 {
   using namespace omp_addr_tokenizer;
   location_t loc = OMP_CLAUSE_LOCATION (c);
   int i = *idx;
   tree decl = addr_tokens[i + 1]->expr;
+  bool decl_p = DECL_P (decl);
   bool declare_target_p = (decl_p
 			   && is_global_var (decl)
 			   && lookup_attribute ("omp declare target",
 						DECL_ATTRIBUTES (decl)));
   bool map_p = OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP;
-  bool implicit_p = (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
-		     && OMP_CLAUSE_MAP_IMPLICIT (c));
+  bool implicit_p = map_p && OMP_CLAUSE_MAP_IMPLICIT (c);
   bool chain_p = omp_access_chain_p (addr_tokens, i + 1);
   tree c2 = NULL_TREE, c3 = NULL_TREE;
   unsigned consume_tokens = 2;
-  bool target = (ort & C_ORT_TARGET) != 0;
-  bool openmp = (ort & C_ORT_OMP) != 0;
+  bool target_p = (ort & C_ORT_TARGET) != 0;
+  bool openmp_p = (ort & C_ORT_OMP) != 0;
 
   gcc_assert (i == 0);
 
-  if (!openmp
-      && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+  if (!openmp_p
+      && map_p
       && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
 	  || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))
     {
-      *idx = ++i;
+      i += 2;
+      *idx = i;
       return c;
     }
 
   switch (addr_tokens[i + 1]->u.access_kind)
     {
     case ACCESS_DIRECT:
-      if (decl_p && !target)
+      if (decl_p && !target_p)
 	c_common_mark_addressable_vec (addr_tokens[i + 1]->expr);
       break;
 
@@ -3573,7 +3585,11 @@ c_omp_address_inspector::expand_array_base (tree c,
 	/* Copy the referenced object.  Note that we do this even for !MAP_P
 	   clauses.  */
 	tree obj = convert_from_reference (addr_tokens[i + 1]->expr);
-	OMP_CLAUSE_DECL (c) = obj;
+	if (TREE_CODE (TREE_TYPE (obj)) == ARRAY_TYPE)
+	  /* We have a ref to array: add a [0] element as the ME expects.  */
+	  OMP_CLAUSE_DECL (c) = build_array_ref (loc, obj, integer_zero_node);
+	else
+	  OMP_CLAUSE_DECL (c) = obj;
 	OMP_CLAUSE_SIZE (c) = TYPE_SIZE_UNIT (TREE_TYPE (obj));
 
 	if (!map_p)
@@ -3583,6 +3599,9 @@ c_omp_address_inspector::expand_array_base (tree c,
 	    break;
 	  }
 
+	if (!target_p)
+	  break;
+
 	/* If we have a reference to a pointer, avoid using
 	   FIRSTPRIVATE_REFERENCE here in case the pointer is modified in the
 	   offload region (we can only do that if the pointer does not point
@@ -3590,32 +3609,28 @@ c_omp_address_inspector::expand_array_base (tree c,
 	   FROM mapping...  */
 	bool ref_to_ptr = TREE_CODE (TREE_TYPE (obj)) == POINTER_TYPE;
 
-	if (target)
+	c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
+	if (!ref_to_ptr
+	    && !declare_target_p
+	    && decl_p)
+	  OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_REFERENCE);
+	else
 	  {
-	    c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
-	    if (target
-		&& !ref_to_ptr
-		&& !declare_target_p
-		&& decl_p)
-	      OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_REFERENCE);
-	    else
-	      {
-		OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
-		if (decl_p)
-		  c_common_mark_addressable_vec (addr_tokens[i + 1]->expr);
-	      }
-	    OMP_CLAUSE_DECL (c2) = addr_tokens[i + 1]->expr;
-	    OMP_CLAUSE_SIZE (c2) = size_zero_node;
+	    OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
+	    if (decl_p)
+	      c_common_mark_addressable_vec (addr_tokens[i + 1]->expr);
+	  }
+	OMP_CLAUSE_DECL (c2) = addr_tokens[i + 1]->expr;
+	OMP_CLAUSE_SIZE (c2) = size_zero_node;
 
-	    if (ref_to_ptr)
-	      {
-		c3 = c2;
-		c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
-		OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALLOC);
-		OMP_CLAUSE_DECL (c2) = addr_tokens[i + 1]->expr;
-		OMP_CLAUSE_SIZE (c2)
-		  = TYPE_SIZE_UNIT (TREE_TYPE (OMP_CLAUSE_DECL (c2)));
-	      }
+	if (ref_to_ptr)
+	  {
+	    c3 = c2;
+	    c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
+	    OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALLOC);
+	    OMP_CLAUSE_DECL (c2) = addr_tokens[i + 1]->expr;
+	    OMP_CLAUSE_SIZE (c2)
+	      = TYPE_SIZE_UNIT (TREE_TYPE (OMP_CLAUSE_DECL (c2)));
 	  }
       }
       break;
@@ -3629,6 +3644,9 @@ c_omp_address_inspector::expand_array_base (tree c,
 	    break;
 	  }
 
+	if (!target_p)
+	  break;
+
 	tree virtual_origin
 	  = convert_from_reference (addr_tokens[i + 1]->expr);
 	virtual_origin = build_fold_addr_expr (virtual_origin);
@@ -3636,8 +3654,21 @@ c_omp_address_inspector::expand_array_base (tree c,
 					   virtual_origin);
 	tree data_addr = omp_accessed_addr (addr_tokens, i + 1, expr);
 	c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
-	if (decl_p && target && !declare_target_p)
-	  OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_REFERENCE);
+	if (decl_p && target_p && !declare_target_p)
+	  {
+	    /* It appears that omp-low.cc mishandles cases where we have a
+	       [reference to an] array of pointers such as:
+
+		 int *arr[N];   (or "int *(&arr)[N] = ...")
+		 #pragma omp target map(arr[a][b:c])
+		 { ... }
+
+	       in such cases chain_p will be true.  For now, fall back to
+	       GOMP_MAP_POINTER.  */
+	    enum gomp_map_kind k = chain_p ? GOMP_MAP_POINTER
+					   : GOMP_MAP_FIRSTPRIVATE_REFERENCE;
+	    OMP_CLAUSE_SET_MAP_KIND (c2, k);
+	  }
 	else
 	  {
 	    if (decl_p)
@@ -3665,7 +3696,7 @@ c_omp_address_inspector::expand_array_base (tree c,
 	/* The code handling "firstprivatize_array_bases" in gimplify.cc is
 	   relevant here.  What do we need to create for arrays at this
 	   stage?  (This condition doesn't feel quite right.  FIXME?)  */
-	if (!target
+	if (!target_p
 	    && (TREE_CODE (TREE_TYPE (addr_tokens[i + 1]->expr))
 		== ARRAY_TYPE))
 	  break;
@@ -3676,8 +3707,13 @@ c_omp_address_inspector::expand_array_base (tree c,
 					   virtual_origin);
 	tree data_addr = omp_accessed_addr (addr_tokens, i + 1, expr);
 	c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
-	if (decl_p && target)
-	  OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
+	if (decl_p && target_p)
+	  {
+	    /* See comment for ACCESS_INDEXED_REF_TO_ARRAY above.  */
+	    enum gomp_map_kind k = chain_p ? GOMP_MAP_POINTER
+					   : GOMP_MAP_FIRSTPRIVATE_POINTER;
+	    OMP_CLAUSE_SET_MAP_KIND (c2, k);
+	  }
 	else
 	  {
 	    if (decl_p)
@@ -3730,7 +3766,7 @@ c_omp_address_inspector::expand_array_base (tree c,
 	   regions (e.g. "acc data" constructs).  It'll be removed anyway in
 	   gimplify.cc, but doing it this way maintains diagnostic
 	   behaviour.  */
-	if (decl_p && (target || !openmp) && !chain_p && !declare_target_p)
+	if (decl_p && (target_p || !openmp_p) && !chain_p && !declare_target_p)
 	  OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
 	else
 	  {
@@ -3785,15 +3821,19 @@ c_omp_address_inspector::expand_array_base (tree c,
 
 	tree data_addr = omp_accessed_addr (addr_tokens, last_access, expr);
 	c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
-	if (decl_p && target && !declare_target_p)
-	  OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_REFERENCE);
+	if (decl_p && target_p && !chain_p && !declare_target_p)
+	  {
+	    OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_REFERENCE);
+	    OMP_CLAUSE_DECL (c2) = addr_tokens[i + 1]->expr;
+	  }
 	else
 	  {
 	    if (decl_p)
 	      c_common_mark_addressable_vec (addr_tokens[i + 1]->expr);
 	    OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
+	    OMP_CLAUSE_DECL (c2)
+	      = convert_from_reference (addr_tokens[i + 1]->expr);
 	  }
-	OMP_CLAUSE_DECL (c2) = addr_tokens[i + 1]->expr;
 	OMP_CLAUSE_SIZE (c2)
 	  = fold_build2_loc (loc, MINUS_EXPR, ptrdiff_type_node,
 			     fold_convert_loc (loc, ptrdiff_type_node,
@@ -3832,7 +3872,7 @@ c_omp_address_inspector::expand_array_base (tree c,
   *idx = i;
 
   if (chain_p && map_p)
-    return omp_expand_access_chain (c, expr, addr_tokens, idx);
+    return omp_expand_access_chain (c, expr, addr_tokens, idx, ort);
 
   return c;
 }
@@ -3843,7 +3883,8 @@ tree
 c_omp_address_inspector::expand_component_selector (tree c,
 						    vec<omp_addr_token *>
 						      &addr_tokens,
-						    tree expr, unsigned *idx)
+						    tree expr, unsigned *idx,
+						    c_omp_region_type ort)
 {
   using namespace omp_addr_tokenizer;
   location_t loc = OMP_CLAUSE_LOCATION (c);
@@ -3973,7 +4014,7 @@ c_omp_address_inspector::expand_component_selector (tree c,
   *idx = i;
 
   if (chain_p && map_p)
-    return omp_expand_access_chain (c, expr, addr_tokens, idx);
+    return omp_expand_access_chain (c, expr, addr_tokens, idx, ort);
 
   return c;
 }
@@ -3998,7 +4039,7 @@ c_omp_address_inspector::expand_map_clause (tree c, tree expr,
 	  && addr_tokens[i]->u.structure_base_kind == BASE_DECL
 	  && addr_tokens[i + 1]->type == ACCESS_METHOD)
 	{
-	  c = expand_array_base (c, addr_tokens, expr, &i, ort, true);
+	  c = expand_array_base (c, addr_tokens, expr, &i, ort);
 	  if (c == error_mark_node)
 	    return error_mark_node;
 	}
@@ -4007,7 +4048,7 @@ c_omp_address_inspector::expand_map_clause (tree c, tree expr,
 	       && addr_tokens[i]->u.structure_base_kind == BASE_ARBITRARY_EXPR
 	       && addr_tokens[i + 1]->type == ACCESS_METHOD)
 	{
-	  c = expand_array_base (c, addr_tokens, expr, &i, ort, false);
+	  c = expand_array_base (c, addr_tokens, expr, &i, ort);
 	  if (c == error_mark_node)
 	    return error_mark_node;
 	}
@@ -4043,7 +4084,7 @@ c_omp_address_inspector::expand_map_clause (tree c, tree expr,
 	       && addr_tokens[i]->type == COMPONENT_SELECTOR
 	       && addr_tokens[i + 1]->type == ACCESS_METHOD)
 	{
-	  c = expand_component_selector (c, addr_tokens, expr, &i);
+	  c = expand_component_selector (c, addr_tokens, expr, &i, ort);
 	  /* We used 'expr', so these must have been the last tokens.  */
 	  gcc_assert (i == length);
 	  if (c == error_mark_node)
diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 891d449b1c0..69d2bee9c16 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -19057,7 +19057,7 @@ c_parser_omp_clause_detach (c_parser *parser, tree list)
 static tree
 c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
 			   const char *where, bool finish_p = true,
-			   bool target = false)
+			   bool target_p = false)
 {
   tree clauses = NULL;
   bool first = true;
@@ -19267,8 +19267,8 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
   c_parser_skip_to_pragma_eol (parser);
 
   if (finish_p)
-    return c_finish_omp_clauses (clauses, target ? C_ORT_ACC_TARGET
-						 : C_ORT_ACC);
+    return c_finish_omp_clauses (clauses, target_p ? C_ORT_ACC_TARGET
+						   : C_ORT_ACC);
 
   return clauses;
 }
@@ -23623,6 +23623,7 @@ c_parser_omp_target_data (location_t loc, c_parser *parser, bool *if_p)
 	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
 	  case GOMP_MAP_ALWAYS_POINTER:
 	  case GOMP_MAP_ATTACH_DETACH:
+	  case GOMP_MAP_ATTACH:
 	    break;
 	  default:
 	    map_seen |= 1;
@@ -23863,7 +23864,8 @@ c_parser_omp_target_exit_data (location_t loc, c_parser *parser,
 
   tree clauses
     = c_parser_omp_all_clauses (parser, OMP_TARGET_EXIT_DATA_CLAUSE_MASK,
-				"#pragma omp target exit data");
+				"#pragma omp target exit data", false);
+  clauses = c_finish_omp_clauses (clauses, C_ORT_OMP_EXIT_DATA);
   c_omp_adjust_map_clauses (clauses, false);
   int map_seen = 0;
   for (tree *pc = &clauses; *pc;)
@@ -24155,7 +24157,9 @@ check_clauses:
 	  case GOMP_MAP_PRESENT_ALLOC:
 	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
 	  case GOMP_MAP_ALWAYS_POINTER:
+	  case GOMP_MAP_POINTER:
 	  case GOMP_MAP_ATTACH_DETACH:
+	  case GOMP_MAP_ATTACH:
 	    break;
 	  default:
 	    error_at (OMP_CLAUSE_LOCATION (*pc),
diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc
index 10ca0a2b343..011a1991d8a 100644
--- a/gcc/c/c-typeck.cc
+++ b/gcc/c/c-typeck.cc
@@ -13891,7 +13891,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 /* Handle array sections for clause C.  */
 
 static bool
-handle_omp_array_sections (tree c, enum c_omp_region_type ort)
+handle_omp_array_sections (tree &c, enum c_omp_region_type ort)
 {
   bool maybe_zero_len = false;
   unsigned int first_non_one = 0;
@@ -14100,9 +14100,16 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
       OMP_CLAUSE_DECL (c) = first;
       if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
 	return false;
-      if (size)
-	size = c_fully_fold (size, false, NULL);
-      OMP_CLAUSE_SIZE (c) = size;
+      /* Don't set OMP_CLAUSE_SIZE for bare attach/detach clauses.  */
+      if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+	  || (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH
+	      && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_DETACH
+	      && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FORCE_DETACH))
+	{
+	  if (size)
+	    size = c_fully_fold (size, false, NULL);
+	  OMP_CLAUSE_SIZE (c) = size;
+	}
 
       if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
 	return false;
@@ -14117,9 +14124,21 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
       tree nc = ai.expand_map_clause (c, first, addr_tokens, ort);
       if (nc != error_mark_node)
 	{
+	  using namespace omp_addr_tokenizer;
+
 	  if (ai.maybe_zero_length_array_section (c))
 	    OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1;
 
+	  /* !!! If we're accessing a base decl via chained access
+	     methods (e.g. multiple indirections), duplicate clause
+	     detection won't work properly.  Skip it in that case.  */
+	  if ((addr_tokens[0]->type == STRUCTURE_BASE
+	       || addr_tokens[0]->type == ARRAY_BASE)
+	      && addr_tokens[0]->u.structure_base_kind == BASE_DECL
+	      && addr_tokens[1]->type == ACCESS_METHOD
+	      && omp_access_chain_p (addr_tokens, 1))
+	    c = nc;
+
 	  return false;
 	}
     }
@@ -15289,7 +15308,8 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		  remove = true;
 		if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 		    && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
-			|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))
+			|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
+		    && !OMP_CLAUSE_SIZE (c))
 		  /* In this case, we have a single array element which is a
 		     pointer, and we already set OMP_CLAUSE_SIZE in
 		     handle_omp_array_sections above.  For attach/detach
@@ -15319,7 +15339,8 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	      }
 	    if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 		&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
-		    || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))
+		    || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
+		&& !OMP_CLAUSE_SIZE (c))
 	      /* For attach/detach clauses, set OMP_CLAUSE_SIZE (representing a
 		 bias) to zero here, so it is not set erroneously to the pointer
 		 size later on in gimplify.cc.  */
@@ -15432,7 +15453,9 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		  bitmap_set_bit (&map_firstprivate_head, DECL_UID (t));
 	      }
 	    else if (bitmap_bit_p (&map_head, DECL_UID (t))
-		     && !bitmap_bit_p (&map_field_head, DECL_UID (t)))
+		     && !bitmap_bit_p (&map_field_head, DECL_UID (t))
+		     && ort != C_ORT_OMP
+		     && ort != C_ORT_OMP_EXIT_DATA)
 	      {
 		if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
 		  error_at (OMP_CLAUSE_LOCATION (c),
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index d25a6c7734a..e54acd39025 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -41358,7 +41358,7 @@ cp_parser_oacc_compute_clause_self (cp_parser *parser, tree list)
 static tree
 cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
 			    const char *where, cp_token *pragma_tok,
-			    bool finish_p = true, bool target = false)
+			    bool finish_p = true, bool target_p = false)
 {
   tree clauses = NULL;
   bool first = true;
@@ -41577,7 +41577,8 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
   cp_parser_skip_to_pragma_eol (parser, pragma_tok);
 
   if (finish_p)
-    return finish_omp_clauses (clauses, target ? C_ORT_ACC_TARGET : C_ORT_ACC);
+    return finish_omp_clauses (clauses, target_p ? C_ORT_ACC_TARGET
+						 : C_ORT_ACC);
 
   return clauses;
 }
@@ -46121,6 +46122,7 @@ cp_parser_omp_target_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p)
 	  case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
 	  case GOMP_MAP_ALWAYS_POINTER:
 	  case GOMP_MAP_ATTACH_DETACH:
+	  case GOMP_MAP_ATTACH:
 	    break;
 	  default:
 	    map_seen |= 1;
@@ -46322,7 +46324,9 @@ cp_parser_omp_target_exit_data (cp_parser *parser, cp_token *pragma_tok,
 
   tree clauses
     = cp_parser_omp_all_clauses (parser, OMP_TARGET_EXIT_DATA_CLAUSE_MASK,
-				 "#pragma omp target exit data", pragma_tok);
+				 "#pragma omp target exit data", pragma_tok,
+				 false);
+  clauses = finish_omp_clauses (clauses, C_ORT_OMP_EXIT_DATA);
   c_omp_adjust_map_clauses (clauses, false);
   int map_seen = 0;
   for (tree *pc = &clauses; *pc;)
diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc
index 3f90fdbe51f..ae85df0590a 100644
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -5903,7 +5903,11 @@ handle_omp_array_sections (tree &c, enum c_omp_region_type ort)
 	  OMP_CLAUSE_DECL (c) = first;
 	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
 	    return false;
-	  OMP_CLAUSE_SIZE (c) = size;
+	  if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+	      || (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH
+		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_DETACH
+		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FORCE_DETACH))
+	    OMP_CLAUSE_SIZE (c) = size;
 	  if (TREE_CODE (t) == FIELD_DECL)
 	    t = finish_non_static_data_member (t, NULL_TREE, NULL_TREE);
 
@@ -8271,7 +8275,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		  remove = true;
 		if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 		    && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
-			|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))
+			|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
+		    && !OMP_CLAUSE_SIZE (c))
 		  /* In this case, we have a single array element which is a
 		     pointer, and we already set OMP_CLAUSE_SIZE in
 		     handle_omp_array_sections above.  For attach/detach
@@ -8303,7 +8308,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	      }
 	    if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 		&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
-		    || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))
+		    || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
+		&& !OMP_CLAUSE_SIZE (c))
 	      /* For attach/detach clauses, set OMP_CLAUSE_SIZE (representing a
 		 bias) to zero here, so it is not set erroneously to the
 		 pointer size later on in gimplify.cc.  */
@@ -8458,7 +8464,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	      bitmap_set_bit (&map_firstprivate_head, DECL_UID (t));
 	    else if (bitmap_bit_p (&map_head, DECL_UID (t))
 		     && !bitmap_bit_p (&map_field_head, DECL_UID (t))
-		     && ort != C_ORT_OMP)
+		     && ort != C_ORT_OMP
+		     && ort != C_ORT_OMP_EXIT_DATA)
 	      {
 		if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
 		  error_at (OMP_CLAUSE_LOCATION (c),
@@ -10114,6 +10121,7 @@ finish_omp_target (location_t loc, tree clauses, tree body, bool combined_p)
 	  case GOMP_MAP_ATTACH_DETACH:
 	  case GOMP_MAP_ATTACH:
 	  case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
+	  case GOMP_MAP_POINTER:
 	  case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION:
 	    break;
 	  default:
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 15ca2832608..1e34bde19be 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -10608,6 +10608,10 @@ omp_resolve_clause_dependencies (enum tree_code code,
 	  }
 	  break;
 
+	case GOMP_MAP_ATTACH:
+	  /* Ignore standalone attach here.  */
+	  break;
+
 	default:
 	  {
 	    omp_mapping_group *struct_group;
@@ -10869,59 +10873,6 @@ omp_siblist_move_concat_nodes_after (tree first_new, tree *last_new_tail,
   return continue_at;
 }
 
-/* Expand a chained access.  We only expect to see a quite limited range of
-   expression types here, because e.g. you can't have an array of
-   references.  See also c-omp.cc:omp_expand_access_chain.  */
-
-static void
-omp_expand_access_chain (location_t loc, tree **list_pp, tree expr,
-			 vec<omp_addr_token *> &addr_tokens,
-			 unsigned *idx, gomp_map_kind kind)
-{
-  using namespace omp_addr_tokenizer;
-  unsigned i = *idx;
-  tree c = NULL_TREE;
-
-  switch (addr_tokens[i]->u.access_kind)
-    {
-    case ACCESS_POINTER:
-    case ACCESS_POINTER_OFFSET:
-      {
-	tree virtual_origin
-	  = fold_convert_loc (loc, ptrdiff_type_node, addr_tokens[i]->expr);
-	tree data_addr = omp_accessed_addr (addr_tokens, i, expr);
-	c = build_omp_clause (loc, OMP_CLAUSE_MAP);
-	OMP_CLAUSE_SET_MAP_KIND (c, kind);
-	OMP_CLAUSE_DECL (c) = addr_tokens[i]->expr;
-	OMP_CLAUSE_SIZE (c)
-	  = fold_build2_loc (loc, MINUS_EXPR, ptrdiff_type_node,
-			     fold_convert_loc (loc, ptrdiff_type_node,
-					       data_addr),
-			     virtual_origin);
-      }
-      break;
-
-    case ACCESS_INDEXED_ARRAY:
-      break;
-
-    default:
-      return;
-    }
-
-  if (c)
-    {
-      OMP_CLAUSE_CHAIN (c) = **list_pp;
-      **list_pp = c;
-      *list_pp = &OMP_CLAUSE_CHAIN (c);
-    }
-
-  *idx = ++i;
-
-  if (addr_tokens[i]->type == ACCESS_METHOD
-      && omp_access_chain_p (addr_tokens, i))
-    omp_expand_access_chain (loc, list_pp, expr, addr_tokens, idx, kind);
-}
-
 static omp_addr_token *
 omp_first_chained_access_token (vec<omp_addr_token *> &addr_tokens)
 {
@@ -12286,8 +12237,19 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		}
 	    }
 	  if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
-	    OMP_CLAUSE_SIZE (c) = DECL_P (decl) ? DECL_SIZE_UNIT (decl)
-				  : TYPE_SIZE_UNIT (TREE_TYPE (decl));
+	    {
+	      /* Sanity check: attach/detach map kinds use the size as a bias,
+		 and it's never right to use the decl size for such
+		 mappings.  */
+	      gcc_assert (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH
+			  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_DETACH
+			  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FORCE_DETACH
+			  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH_DETACH
+			  && (OMP_CLAUSE_MAP_KIND (c)
+			      != GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION));
+	      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)
 	    {
@@ -12554,6 +12516,32 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 					     ? GOMP_MAP_DETACH
 					     : GOMP_MAP_ATTACH);
 	      OMP_CLAUSE_SET_MAP_KIND (c, map_kind);
+
+	      /* If we have attach/detach but the decl we have is a pointer to
+		 pointer, we're probably mapping the "base level" array
+		 implicitly.  Make sure we don't add the decl as if we mapped
+		 it explicitly.  That is,
+
+		   int **arr;
+		   [...]
+		   #pragma omp target map(arr[a][b:c])
+
+		 should *not* map "arr" explicitly.  That way we get a
+		 zero-length "alloc" mapping for it, and assuming it's been
+		 mapped by some previous directive, etc., things work as they
+		 should.  */
+
+	      tree basetype = TREE_TYPE (addr_tokens[0]->expr);
+
+	      if (TREE_CODE (basetype) == REFERENCE_TYPE)
+		basetype = TREE_TYPE (basetype);
+
+	      if (code == OMP_TARGET
+		  && addr_tokens[0]->type == ARRAY_BASE
+		  && addr_tokens[0]->u.structure_base_kind == BASE_DECL
+		  && TREE_CODE (basetype) == POINTER_TYPE
+		  && TREE_CODE (TREE_TYPE (basetype)) == POINTER_TYPE)
+		break;
 	    }
 	  else if ((code == OACC_ENTER_DATA
 		    || code == OACC_EXIT_DATA
@@ -13480,7 +13468,16 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
       OMP_CLAUSE_SIZE (nc) = size_zero_node;
       OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_ALLOC);
       OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (clause) = 1;
-      OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_FIRSTPRIVATE_POINTER);
+      tree dtype = TREE_TYPE (decl);
+      if (TREE_CODE (dtype) == REFERENCE_TYPE)
+	dtype = TREE_TYPE (dtype);
+      /* FIRSTPRIVATE_POINTER doesn't work well if we have a
+	 multiply-indirected pointer.  */
+      if (TREE_CODE (dtype) == POINTER_TYPE
+	  && TREE_CODE (TREE_TYPE (dtype)) == POINTER_TYPE)
+	OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_POINTER);
+      else
+	OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_FIRSTPRIVATE_POINTER);
       OMP_CLAUSE_CHAIN (nc) = chain;
       OMP_CLAUSE_CHAIN (clause) = nc;
       struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
diff --git a/gcc/testsuite/c-c++-common/gomp/target-enter-data-1.c b/gcc/testsuite/c-c++-common/gomp/target-enter-data-1.c
index 37c8dd4e328..4913d338e5f 100644
--- a/gcc/testsuite/c-c++-common/gomp/target-enter-data-1.c
+++ b/gcc/testsuite/c-c++-common/gomp/target-enter-data-1.c
@@ -21,5 +21,5 @@ void func (struct foo *f, int n, int m)
   #pragma omp target enter data map (to: f->bars[n].vectors[:f->bars[n].num_vectors])
 }
 
-/* { dg-final { scan-tree-dump-times {map\(struct:\*f \[len: 1\]\) map\(alloc:[a-z0-9\._]+->vectors \[len: 0\]\) map\(to:\*_[0-9]+ \[len: _[0-9]+\]\) map\(attach:[a-z0-9\._]+->vectors \[bias: [^\]]+\]\) map\(attach:\*_[0-9]+ \[bias: 0\]\)} 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times {map\(struct:\*f \[len: 1\]\) map\(alloc:[a-z0-9\._]+->vectors \[len: 0\]\) map\(to:\*_[0-9]+ \[len: _[0-9]+\]\) map\(attach:[a-z0-9\._]+->vectors \[bias: [^\]]+\]\) map\(attach:\*_[0-9]+ \[bias: _[0-9]+\]\)} 1 "gimple" } } */
 /* { dg-final { scan-tree-dump-times {map\(struct:\*\(f->bars \+ \(sizetype\) \(\([^\)]+\) n \* 16\)\) \[len: 1\]\) map\(alloc:[a-z0-9\._]+->vectors \[len: 0\]\) map\(to:\*_[0-9]+ \[len: _[0-9]+\]\) map\(attach:[a-z0-9\._]+->vectors \[bias: [^\]]+\]\)} 2 "gimple" } } */
diff --git a/libgomp/testsuite/libgomp.c++/baseptrs-8.C b/libgomp/testsuite/libgomp.c++/baseptrs-8.C
new file mode 100644
index 00000000000..f9991818551
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/baseptrs-8.C
@@ -0,0 +1,70 @@
+/* { dg-do run } */
+
+#include <stdlib.h>
+#include <assert.h>
+
+#define N 1024
+#define M 64
+
+int main (void)
+{
+  int *a_orig[N];
+  int *(&a)[N] = a_orig;
+
+  for (int i = 0; i < N; i++)
+    a[i] = (int *) calloc (M, sizeof (int));
+
+  /* 'target enter data'/'target exit data' with array of pointers.  */
+#pragma omp target enter data map(alloc: a[0:N])
+
+  for (int i = 0; i < N; i++)
+  {
+#pragma omp target enter data map(to: a[i][0:M])
+  }
+
+#pragma omp target map(alloc: a)
+  {
+    for (int i = 0; i < N; i++)
+      for (int j = 0; j < M; j++)
+	a[i][j] = i + j;
+  }
+
+for (int i = 0; i < N; i++)
+  {
+#pragma omp target exit data map(release: a[i]) map(from: a[i][0:M])
+  }
+
+#pragma omp target exit data map(release: a, a[0:N])
+
+  /* 'target data' with array of pointers.  */
+#pragma omp target data map(alloc: a[0:N])
+  {
+#pragma omp target data map(tofrom: a[5][0:M])
+    {
+#pragma omp target map(alloc: a)
+      {
+	for (int i = 0; i < M; i++)
+	  a[5][i]++;
+      }
+    }
+  }
+
+  /* 'target' with array of pointers.  */
+#pragma omp target data map(alloc: a[0:N])
+  {
+#pragma omp target map(tofrom: a[7][0:M])
+    {
+      for (int i = 0; i < M; i++)
+	a[7][i] += 2;
+    }
+  }
+
+  for (int i = 0; i < N; i++)
+    for (int j = 0; j < M; j++)
+      assert (a[i][j] == i + j + (i == 5) + 2 * (i == 7));
+
+  for (int i = 0; i < N; i++)
+    free (a[i]);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/baseptrs-9.C b/libgomp/testsuite/libgomp.c++/baseptrs-9.C
new file mode 100644
index 00000000000..95e7eebb0ed
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/baseptrs-9.C
@@ -0,0 +1,57 @@
+/* { dg-do run } */
+
+#include <stdlib.h>
+#include <assert.h>
+
+int main (void)
+{
+  int **a_orig,i,j,n;
+  int **&a = a_orig;
+
+  j = 3;
+  n = 12;
+
+  a = (int **) calloc (32, sizeof (int *));
+  for (int x = 0; x < 32; x++)
+    a[x] = (int *) calloc (32, sizeof (int));
+
+  for (int i = 2; i < 32; i++)
+    {
+      #pragma omp target enter data map(a, a[2:30])
+      #pragma omp target enter data map(a[i][j:n])
+      #pragma omp target map(alloc: a)
+      {
+	for (int x = j; x < j + n; x++)
+	  a[i][x]++;
+      }
+      #pragma omp target exit data map(a[i][j:n])
+
+      #pragma omp target data map(a, a[i][j:n])
+      {
+	#pragma omp target map(alloc: a)
+	{
+	  for (int x = j; x < j + n; x++)
+	    a[i][x]++;
+	}
+      }
+      #pragma omp target exit data map(a, a[2:30])
+
+      #pragma omp target data map(a, a[2:30])
+      {
+	#pragma omp target map(a[i][j:n])
+	{
+	  for (int x = j; x < j + n; x++)
+	    a[i][x]++;
+	}
+      }
+    }
+
+  for (int i = 0; i < 32; i++)
+    for (int j = 0; j < 32; j++)
+      if (i >= 2 && j >= 3 && j < 15)
+	assert (a[i][j] == 3);
+      else
+	assert (a[i][j] == 0);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/ref-mapping-1.C b/libgomp/testsuite/libgomp.c++/ref-mapping-1.C
new file mode 100644
index 00000000000..9aa232f3f67
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/ref-mapping-1.C
@@ -0,0 +1,80 @@
+/* { dg-do run } */
+
+#include <cassert>
+
+void test_ref ()
+{
+  int c_orig = 5;
+  int &c = c_orig;
+
+#pragma omp target map(tofrom: c)
+  {
+    c++;
+  }
+
+  assert (c == 6);
+}
+
+void test_ref_to_ptr ()
+{
+  int val = 5;
+  int *ptr_orig = &val;
+  int *&ptr_ref = ptr_orig;
+
+#pragma omp target map(tofrom: ptr_ref[0])
+  {
+    (*ptr_ref)++;
+  }
+
+  assert (val == 6);
+}
+
+void test_ref_to_array ()
+{
+  int arr[1];
+  int (&arr_ref)[1] = arr;
+
+  arr_ref[0] = 5;
+
+#pragma omp target map(tofrom: arr_ref[0:1])
+  {
+    arr_ref[0]++;
+  }
+
+  assert (arr_ref[0] == 6);
+
+#pragma omp target map(tofrom: arr_ref[0])
+  {
+    arr_ref[0]++;
+  }
+
+  assert (arr_ref[0] == 7);
+}
+
+void test_ref_to_ptr_array ()
+{
+  int *arr[1];
+  int *(&arr_ref)[1] = arr;
+  int val = 5;
+
+  arr_ref[0] = &val;
+
+#pragma omp target data map(alloc: arr_ref, arr_ref[0])
+  {
+#pragma omp target map(tofrom: arr_ref[0][0:1])
+    {
+      arr_ref[0][0]++;
+    }
+  }
+
+  assert (arr_ref[0][0] == 6);
+}
+
+int main ()
+{
+  test_ref ();
+  test_ref_to_ptr ();
+  test_ref_to_array ();
+  test_ref_to_ptr_array ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/baseptrs-6.c b/libgomp/testsuite/libgomp.c-c++-common/baseptrs-6.c
new file mode 100644
index 00000000000..4b6e237471a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/baseptrs-6.c
@@ -0,0 +1,69 @@
+/* { dg-do run } */
+
+#include <stdlib.h>
+#include <assert.h>
+
+#define N 1024
+#define M 64
+
+int main (void)
+{
+  int *a[N];
+
+  for (int i = 0; i < N; i++)
+    a[i] = (int *) calloc (M, sizeof (int));
+
+  /* 'target enter data'/'target exit data' with array of pointers.  */
+#pragma omp target enter data map(alloc: a[0:N])
+
+  for (int i = 0; i < N; i++)
+  {
+#pragma omp target enter data map(to: a[i][0:M])
+  }
+
+#pragma omp target map(alloc: a)
+  {
+    for (int i = 0; i < N; i++)
+      for (int j = 0; j < M; j++)
+	a[i][j] = i + j;
+  }
+
+for (int i = 0; i < N; i++)
+  {
+#pragma omp target exit data map(release: a[i]) map(from: a[i][0:M])
+  }
+
+#pragma omp target exit data map(release: a, a[0:N])
+
+  /* 'target data' with array of pointers.  */
+#pragma omp target data map(alloc: a[0:N])
+  {
+#pragma omp target data map(tofrom: a[5][0:M])
+    {
+#pragma omp target map(alloc: a)
+      {
+	for (int i = 0; i < M; i++)
+	  a[5][i]++;
+      }
+    }
+  }
+
+  /* 'target' with array of pointers.  */
+#pragma omp target data map(alloc: a[0:N])
+  {
+#pragma omp target map(tofrom: a[7][0:M])
+    {
+      for (int i = 0; i < M; i++)
+	a[7][i] += 2;
+    }
+  }
+
+  for (int i = 0; i < N; i++)
+    for (int j = 0; j < M; j++)
+      assert (a[i][j] == i + j + (i == 5) + 2 * (i == 7));
+
+  for (int i = 0; i < N; i++)
+    free (a[i]);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/baseptrs-7.c b/libgomp/testsuite/libgomp.c-c++-common/baseptrs-7.c
new file mode 100644
index 00000000000..9c6710e4e5b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/baseptrs-7.c
@@ -0,0 +1,56 @@
+/* { dg-do run } */
+
+#include <stdlib.h>
+#include <assert.h>
+
+int main (void)
+{
+  int **a,i,j,n;
+
+  j = 3;
+  n = 12;
+
+  a = (int **) calloc (32, sizeof (int *));
+  for (int x = 0; x < 32; x++)
+    a[x] = (int *) calloc (32, sizeof (int));
+
+  for (int i = 2; i < 32; i++)
+    {
+      #pragma omp target enter data map(a, a[2:30])
+      #pragma omp target enter data map(a[i][j:n])
+      #pragma omp target map(alloc: a)
+      {
+	for (int x = j; x < j + n; x++)
+	  a[i][x]++;
+      }
+      #pragma omp target exit data map(a[i][j:n])
+
+      #pragma omp target data map(a, a[i][j:n])
+      {
+	#pragma omp target map(alloc: a)
+	{
+	  for (int x = j; x < j + n; x++)
+	    a[i][x]++;
+	}
+      }
+      #pragma omp target exit data map(a, a[2:30])
+
+      #pragma omp target data map(a, a[0:32])
+      {
+	#pragma omp target map(a[i][j:n])
+	{
+	  for (int x = j; x < j + n; x++)
+	    a[i][x]++;
+	}
+      }
+    }
+
+  for (int i = 0; i < 32; i++)
+    for (int j = 0; j < 32; j++)
+      if (i >= 2 && j >= 3 && j < 15)
+	assert (a[i][j] == 3);
+      else
+	assert (a[i][j] == 0);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.fortran/target-enter-data-6.f90 b/libgomp/testsuite/libgomp.fortran/target-enter-data-6.f90
index 80d30edbfc5..b55d0b268d4 100644
--- a/libgomp/testsuite/libgomp.fortran/target-enter-data-6.f90
+++ b/libgomp/testsuite/libgomp.fortran/target-enter-data-6.f90
@@ -3,6 +3,16 @@
 ! - arrays with array descriptors
 ! For those, the array descriptor / string length must be mapped with 'to:'
 
+! This test fails without the following additional patches:
+!
+! "OpenMP: Pointers and member mappings":
+! https://gcc.gnu.org/pipermail/gcc-patches/2023-August/627898.html
+!
+! "OpenMP/OpenACC: Reorganise OMP map clause handling in gimplify.cc":
+! https://gcc.gnu.org/pipermail/gcc-patches/2023-August/627900.html
+!
+! { dg-xfail-run-if TODO { offload_device_nonshared_as } }
+
 program main
 implicit none