OpenMP: Handle same-directive mapped vars with pointer predefined firstprivate [PR110639]

Predefined 'firstprivate' for pointer variables firstprivatizes the pointer
but if it is associated with a mapped target, its address is updated to the
corresponding target. (If not, the host value remains.)

This commit extends this handling to also update the pointer address for
storaged mapped on the same directive.

The 'gimplify_scan_omp_clauses' change avoids adding an additional
  map(alloc:this) (+ptr assignment)
when there is already a
  map(tofrom:*this) (+ptr assignment)
This shows up for libgomp.c++/pr108286.C and also when 'this' is
actually '__closure->this' (-> g++.dg/gomp/target-{this-{2,4},lambda-1}.C).

	PR middle-end/110639

gcc/ChangeLog:

	* gimplify.cc (struct gimplify_adjust_omp_clauses_data): Add
	append_list.
	(gimplify_adjust_omp_clauses_1, gimplify_adjust_omp_clauses): Add
	GOVD_MAP_0LEN_ARRAY clauses at the end.
	(gimplify_scan_omp_clauses): Mark also '*var' as found not only
	'var'.

libgomp/ChangeLog:

	* target.c (gomp_map_vars_internal): Handle also variables
	mapped in the same directive for GOVD_MAP_0LEN_ARRAY.
	* testsuite/libgomp.c++/pr108286.C: Add gimple tree-scan test.
	* testsuite/libgomp.c-c++-common/target-implicit-map-6.c: New test.

gcc/testsuite/ChangeLog:

        * g++.dg/gomp/target-this-2.C: Remove 'this' pointer mapping alreay
	mapped via __closure->this.
        * g++.dg/gomp/target-this-4.C: Likewise.
        * g++.dg/gomp/target-lambda-1.C: Likewise. Move 'iptr' pointer
	mapping to the end in scan-tree-dump.

 gcc/gimplify.cc                                    |  45 ++++-
 gcc/testsuite/g++.dg/gomp/target-lambda-1.C        |   4 +-
 gcc/testsuite/g++.dg/gomp/target-this-2.C          |   4 +-
 gcc/testsuite/g++.dg/gomp/target-this-4.C          |   6 +-
 libgomp/target.c                                   |  11 +-
 libgomp/testsuite/libgomp.c++/pr108286.C           |   4 +
 .../libgomp.c-c++-common/target-implicit-map-6.c   | 212 +++++++++++++++++++++
 7 files changed, 276 insertions(+), 10 deletions(-)

diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 342e43a7f25..2234fd6b7e1 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -11586,6 +11586,23 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	  else if (!DECL_P (decl))
 	    {
 	      tree d = decl, *pd;
+	      pd = &OMP_CLAUSE_DECL (c);
+	      if (TREE_CODE (decl) == INDIRECT_REF)
+		{
+		  tree d2 = TREE_OPERAND (decl, 0);
+		  STRIP_NOPS (d2);
+		  if (DECL_P (d2))
+		    {
+		      if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue,
+					 fb_lvalue) == GS_ERROR)
+			{
+			  remove = true;
+			  break;
+			}
+		      decl = d2;
+		      goto handle_map_decl;
+		    }
+		}
 	      if (TREE_CODE (d) == ARRAY_REF)
 		{
 		  while (TREE_CODE (d) == ARRAY_REF)
@@ -11594,7 +11611,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		      && TREE_CODE (TREE_TYPE (d)) == ARRAY_TYPE)
 		    decl = d;
 		}
-	      pd = &OMP_CLAUSE_DECL (c);
 	      if (d == decl
 		  && TREE_CODE (decl) == INDIRECT_REF
 		  && TREE_CODE (TREE_OPERAND (decl, 0)) == COMPONENT_REF
@@ -11774,6 +11790,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		}
 	      break;
 	    }
+	handle_map_decl:
 	  flags = GOVD_MAP | GOVD_EXPLICIT;
 	  if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TO
 	      || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TOFROM)
@@ -11806,7 +11823,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	      OMP_CLAUSE_SET_MAP_KIND (c, map_kind);
 	    }
 
-	  goto do_add;
+	  goto do_add_decl;
 
 	case OMP_CLAUSE_AFFINITY:
 	  gimplify_omp_affinity (list_p, pre_p);
@@ -12571,6 +12588,7 @@ omp_find_stores_stmt (gimple_stmt_iterator *gsi_p,
 struct gimplify_adjust_omp_clauses_data
 {
   tree *list_p;
+  tree append_list;
   gimple_seq *pre_p;
 };
 
@@ -12691,6 +12709,7 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
       && omp_shared_to_firstprivate_optimizable_decl_p (decl))
     omp_mark_stores (gimplify_omp_ctxp->outer_context, decl);
 
+  bool len0_append_list_used = false;
   tree chain = *list_p;
   clause = build_omp_clause (input_location, code);
   OMP_CLAUSE_DECL (clause) = decl;
@@ -12707,6 +12726,11 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
     OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (clause) = 1;
   else if (code == OMP_CLAUSE_MAP && (flags & GOVD_MAP_0LEN_ARRAY) != 0)
     {
+      /* For GOVD_MAP_0LEN_ARRAY, add the clauses to append_list such
+	 that those come after any data mapping.  */
+      len0_append_list_used = true;
+      struct gimplify_adjust_omp_clauses_data *adjdata
+	= (struct gimplify_adjust_omp_clauses_data *) data;
       tree nc = build_omp_clause (input_location, OMP_CLAUSE_MAP);
       OMP_CLAUSE_DECL (nc) = decl;
       if (TREE_CODE (TREE_TYPE (decl)) == REFERENCE_TYPE
@@ -12721,8 +12745,9 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
       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);
-      OMP_CLAUSE_CHAIN (nc) = chain;
+      OMP_CLAUSE_CHAIN (nc) = adjdata->append_list;
       OMP_CLAUSE_CHAIN (clause) = nc;
+      adjdata->append_list = clause;
       struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
       gimplify_omp_ctxp = ctx->outer_context;
       gimplify_expr (&TREE_OPERAND (OMP_CLAUSE_DECL (clause), 0),
@@ -12833,7 +12858,8 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
 					  (ctx->region_type & ORT_ACC) != 0);
       gimplify_omp_ctxp = ctx;
     }
-  *list_p = clause;
+  if (!len0_append_list_used)
+    *list_p = clause;
   struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
   gimplify_omp_ctxp = ctx->outer_context;
   /* Don't call omp_finish_clause on implicitly added OMP_CLAUSE_PRIVATE
@@ -12842,7 +12868,7 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
   if (code != OMP_CLAUSE_PRIVATE || ctx->region_type != ORT_SIMD) 
     lang_hooks.decls.omp_finish_clause (clause, pre_p,
 					(ctx->region_type & ORT_ACC) != 0);
-  if (gimplify_omp_ctxp)
+  if (gimplify_omp_ctxp && !len0_append_list_used)
     for (; clause != chain; clause = OMP_CLAUSE_CHAIN (clause))
       if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP
 	  && DECL_P (OMP_CLAUSE_SIZE (clause)))
@@ -13445,6 +13471,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
 
   /* Add in any implicit data sharing.  */
   struct gimplify_adjust_omp_clauses_data data;
+  data.append_list = NULL;
   if ((gimplify_omp_ctxp->region_type & ORT_ACC) == 0)
     {
       /* OpenMP.  Implicit clauses are added at the start of the clause list,
@@ -13472,6 +13499,14 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
 		    "iterator");
 	  break;
 	}
+  if (data.append_list != NULL_TREE && *data.list_p != NULL_TREE)
+    {
+      for (c = *data.list_p; c && OMP_CLAUSE_CHAIN (c); c = OMP_CLAUSE_CHAIN (c))
+	;
+      OMP_CLAUSE_CHAIN (c) = data.append_list;
+    }
+  else if (data.append_list != NULL_TREE)
+    *data.list_p = data.append_list;
 
   gimplify_omp_ctxp = ctx->outer_context;
   delete_omp_context (ctx);
diff --git a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
index 5ce8ceadb19..b4f1593af02 100644
--- a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
+++ b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
@@ -87,7 +87,9 @@ int main (void)
   return 0;
 }
 
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(b\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:iptr \[pointer assign, bias: 0\]\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:__closure->__iptr \[bias: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\)} "gimple" } } */
+/* Note that 'this' = '__closure->__this' such that no pointer-assign for 'this' should appear.  */
+
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(b\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:__closure->__iptr \[bias: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:iptr \[pointer assign, bias: 0\]\)[\r\n]} "gimple" } } */
 
 /* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(end\) firstprivate\(begin\) map\(to:loop \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:loop\.__data1 \[bias: 0\]\)} "gimple" } } */
 
diff --git a/gcc/testsuite/g++.dg/gomp/target-this-2.C b/gcc/testsuite/g++.dg/gomp/target-this-2.C
index cc08e7e8693..eecab5a25e8 100644
--- a/gcc/testsuite/g++.dg/gomp/target-this-2.C
+++ b/gcc/testsuite/g++.dg/gomp/target-this-2.C
@@ -46,4 +46,6 @@ int main (void)
   return 0;
 }
 
-/* { dg-final { scan-tree-dump {map\(alloc:MEM\[\(char \*\)_[0-9]+\] \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) firstprivate\(m\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(from:v \[len: [0-9]+\]\)} "gimple" } } */
+/* Note that 'this' = '__closure->__this' such that no pointer-assign for 'this' should appear.  */
+
+/* { dg-final { scan-tree-dump {firstprivate\(n\) firstprivate\(m\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(from:v \[len: [0-9]+\]\)[\r\n]} "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/target-this-4.C b/gcc/testsuite/g++.dg/gomp/target-this-4.C
index 9ade3cc0b2b..845f89f6997 100644
--- a/gcc/testsuite/g++.dg/gomp/target-this-4.C
+++ b/gcc/testsuite/g++.dg/gomp/target-this-4.C
@@ -102,6 +102,8 @@ int main (void)
   return 0;
 }
 
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(from:mapped \[len: 1\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\)} "gimple" } } */
+/* Note that 'this' = '__closure->__this' such that no pointer-assign for 'this' should appear.  */
 
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:_[0-9]+->refptr \[bias: 0\]\)} "gimple" } } */
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(from:mapped \[len: 1\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\)[\r\n]} "gimple" } } */
+
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:_[0-9]+->refptr \[bias: 0\]\)[\r\n]} "gimple" } } */
diff --git a/libgomp/target.c b/libgomp/target.c
index f30c20255d3..c1f26e6860b 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -1149,7 +1149,8 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
       splay_tree_key n;
       if ((kind & typemask) == GOMP_MAP_ZERO_LEN_ARRAY_SECTION)
 	{
-	  n = gomp_map_0len_lookup (mem_map, &cur_node);
+	  /* Defer lookup when mapped item found.  */
+	  n = not_found_cnt ? NULL : gomp_map_0len_lookup (mem_map, &cur_node);
 	  if (!n)
 	    {
 	      tgt->list[i].key = NULL;
@@ -1417,7 +1418,15 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
 		  }
 		continue;
 	      case GOMP_MAP_FIRSTPRIVATE_INT:
+		continue;
 	      case GOMP_MAP_ZERO_LEN_ARRAY_SECTION:
+		cur_node.host_start = (uintptr_t) hostaddrs[i];
+		cur_node.host_end = cur_node.host_start + sizes[i];
+		n = gomp_map_0len_lookup (mem_map, &cur_node);
+		if (n)
+		  gomp_map_vars_existing (devicep, aq, n, &cur_node,
+					  &tgt->list[i], kind & typemask, false,
+					  implicit, NULL, refcount_set);
 		continue;
 	      case GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT:
 		/* The OpenACC 'host_data' construct only allows 'use_device'
diff --git a/libgomp/testsuite/libgomp.c++/pr108286.C b/libgomp/testsuite/libgomp.c++/pr108286.C
index ee88c2f9fd0..3ae2f5dd16e 100644
--- a/libgomp/testsuite/libgomp.c++/pr108286.C
+++ b/libgomp/testsuite/libgomp.c++/pr108286.C
@@ -1,5 +1,6 @@
 // PR c++/108286
 // { dg-do run }
+// { dg-additional-options "-fdump-tree-gimple" }
 
 struct S {
   int
@@ -27,3 +28,6 @@ main ()
   if (s.foo () != 42)
     __builtin_abort ();
 }
+
+/* Ensure that 'this' is mapped but only once and not additionally via 'this[:0]'.  */
+/* { dg-final { scan-tree-dump "#pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) map\\(tofrom:\\*this \\\[len: \[0-9\]+\\\]\\) map\\(firstprivate:this \\\[pointer assign, bias: 0\\\]\\) nowait map\\(tofrom:res \\\[len: \[0-9\]+\\\]\\) map\\(tofrom:\\*_\[0-9\]+ \\\[len: _\[0-9\]+\\\]\\) map\\(attach:this->ptr \\\[bias: 0\\\]\\)\[\r\n\]" "gimple" } } */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-6.c b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-6.c
new file mode 100644
index 00000000000..6a2d3bfd0e3
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-6.c
@@ -0,0 +1,212 @@
+/* Prefined firstprivate privatizes the pointer
+   and then updates the value to point to the corresponding
+   device variable, if existing.
+
+   See PR middle-end/110639
+   and TR12 in "14.8 target Construct" [379:8-10]
+   or  OpenMP 5.1 in "2.21.7.2 Pointer Initialization for Device Data Environments".  */
+
+#include <stdlib.h>
+#include <omp.h>
+
+int my_false = 0;
+
+int
+f (int x, int y)
+{
+  return x + y;
+}
+
+void
+no_other_clause ()
+{
+  int data = 3;
+  int *p = &data;
+  #pragma omp target enter data map(data)
+  #pragma omp target
+   *p = 5;
+  #pragma omp target exit data map(data)
+  if (*p != 5)
+    abort ();
+}
+
+void
+test1 (int devnum)
+{
+  int start = 0, n = 100;
+  int a[100];
+  int *p = &a[0];
+
+  for (int i = start; i < start+n; i++)
+    a[i] = 10*i;
+
+  #pragma omp target map(a) device(device_num : devnum)
+  {
+    if (my_false)  /* Ensure that 'map(a)' is not optimized away. */
+      a[8] = 1;
+    for (int i = start; i < start+n; i++)
+      p[i] = f(p[i], i);
+    p = NULL;
+  }
+
+  if (p != &a[0])
+    abort ();
+  for (int i = start; i < start+n; i++)
+    if (a[i] != f(10 *i, i))
+      abort ();
+}
+
+
+
+void
+test2 (int devnum)
+{
+  int start = 0, n = 100;
+  int a[100];
+  int *p = &a[0];
+
+  for (int i = start; i < start+n; i++)
+    a[i] = 10*i;
+
+  #pragma omp target enter data map(a) device(device_num : devnum)
+  #pragma omp target device(device_num : devnum)
+  {
+    (void) a; /* Ensure that 'map(a)' is not optimized away. */
+    for (int i = start; i < start+n; i++)
+      p[i] = f(p[i], i);
+    p = NULL;
+  }
+  #pragma omp target exit data map(a) device(device_num : devnum)
+
+  if (p != &a[0])
+    abort ();
+  for (int i = start; i < start+n; i++)
+    if (a[i] != f(10 *i, i))
+      abort ();
+}
+
+void
+test3 (int devnum)
+{
+  int start = 8, n = 10;
+  int a[100];
+  int *p = &a[start];
+
+  for (int i = start; i < start+n; i++)
+    a[i] = 10*i;
+
+  /* p points to a[start]  */
+  #pragma omp target map(a[start:n]) device(device_num : devnum)
+  {
+    if (my_false)  /* Ensure that 'map(a)' is not optimized away. */
+      a[8] = 1;
+    for (int i = 0; i < n; i++)
+      p[i] = f(p[i], i + start);
+    p = NULL;
+  }
+
+  if (p != &a[start])
+    abort ();
+  for (int i = start; i < start+n; i++)
+    if (a[i] != f(10 *i, i))
+      abort ();
+}
+
+void
+test4 (int devnum)
+{
+  int start = 8, n = 10;
+  int a[100];
+  int *p = &a[start];
+
+  for (int i = start; i < start+n; i++)
+    a[i] = 10*i;
+
+  /* p points to a[start]  */
+  #pragma omp target enter data map(a[start:n]) device(device_num : devnum)
+  #pragma omp target device(device_num : devnum)
+  {
+    for (int i = 0; i < n; i++)
+      p[i] = f(p[i], i + start);
+    p = NULL;
+  }
+  #pragma omp target exit data map(a[start:n]) device(device_num : devnum)
+
+  if (p != &a[start])
+    abort ();
+  for (int i = start; i < start+n; i++)
+    if (a[i] != f(10 *i, i))
+      abort ();
+}
+
+void
+test5 (int devnum)
+{
+  int start = 8, n = 10;
+  int a[100];
+  int *p = &a[start + 5];
+
+  for (int i = start; i < start+n; i++)
+    a[i] = 10*i;
+
+  /* p points to a[start + 5]  */
+  #pragma omp target map(a[start:n]) device(device_num : devnum)
+  {
+    if (my_false)  /* Ensure that 'map(a)' is not optimized away. */
+      a[8] = 1;
+    for (int i = 0; i < n; i++)
+      p[i - 5] = f(p[i - 5], i + start);
+    p = NULL;
+  }
+
+  if (p != &a[start + 5])
+    abort ();
+  for (int i = start; i < start+n; i++)
+    if (a[i] != f(10 *i, i))
+      abort ();
+}
+
+void
+test6 (int devnum)
+{
+  int start = 8, n = 10;
+  int a[100];
+  int *p = &a[start + 5];
+
+  for (int i = start; i < start+n; i++)
+    a[i] = 10*i;
+
+  /* p points to a[start + 5]  */
+  #pragma omp target enter data map(a[start:n]) device(device_num : devnum)
+  #pragma omp target device(device_num : devnum)
+  {
+    for (int i = 0; i < n; i++)
+      p[i - 5] = f(p[i - 5], i + start);
+    p = NULL;
+  }
+  #pragma omp target exit data map(a[start:n]) device(device_num : devnum)
+
+  if (p != &a[start + 5])
+    abort ();
+  for (int i = start; i < start+n; i++)
+    if (a[i] != f(10 *i, i))
+      abort ();
+}
+
+int
+main ()
+{
+  int n = omp_get_num_devices ();
+  no_other_clause ();
+  for (int i = omp_initial_device; i <= n; i++)
+    {
+      /* First with 'a' mapped on target; then 'a' on target enter data.  */
+      test1 (i); /* p = &a[0], map(a) */
+      test2 (i);
+      test3 (i); /* p = &a[start], map(a[start:n]) */
+      test4 (i);
+      test5 (i); /* p = &a[start + 5], map(a[start:n]) */
+      test6 (i);
+    }
+  return 0;
+}
