OpenMP 5.0: Remove array section base-pointer mapping semantics and other front-end adjustments

This patch implements three pieces of functionality:

(1) Adjust array section mapping to have standards conforming behavior,
mapping array sections should *NOT* also map the base-pointer:

struct S { int *ptr; ... };
struct S s;

Instead of generating this during gimplify:
                              map(to:*_1 [len: 400]) map(attach:s.ptr [bias: 0])

Now, adjust to:

(i.e. do not map the base-pointer together. The attach operation is still
generated, and if s.ptr is already mapped prior, attachment will happen)

The correct way of achieving the base-pointer-also-mapped behavior would be to
use:

(A small Fortran front-end patch to trans-openmp.c:gfc_trans_omp_array_section
 is also included, which removes generation of a GOMP_MAP_ALWAYS_POINTER for
 array types, which appears incorrect and causes a regression in
 libgomp.fortranlibgomp.fortran/struct-elem-map-1.f90)

(2) Related to the first item above, are fixes in libgomp/target.c to not
overwrite attached pointers when handling device<->host copies, mainly for the
"always" case.

(3) The third is a set of changes to the C/C++ front-ends to extend the allowed
component access syntax in map clauses. These changes are enabled for both
OpenACC and OpenMP.

gcc/c/ChangeLog:

	* c-parser.c (struct omp_dim): New struct type for use inside
	c_parser_omp_variable_list.
	(c_parser_omp_variable_list): Allow multiple levels of array and
	component accesses in array section base-pointer expression.
	(c_parser_omp_clause_to): Set 'allow_deref' to true in call to
	c_parser_omp_var_list_parens.
	(c_parser_omp_clause_from): Likewise.
	* c-typeck.c (handle_omp_array_sections_1): Extend allowed range
	of base-pointer expressions involving INDIRECT/MEM/ARRAY_REF and
	POINTER_PLUS_EXPR.
	(c_finish_omp_clauses): Extend allowed ranged of expressions
	involving INDIRECT/MEM/ARRAY_REF and POINTER_PLUS_EXPR.

gcc/cp/ChangeLog:

	* parser.c (struct omp_dim): New struct type for use inside
	cp_parser_omp_var_list_no_open.
	(cp_parser_omp_var_list_no_open): Allow multiple levels of array and
	component accesses in array section base-pointer expression.
	(cp_parser_omp_all_clauses): Set 'allow_deref' to true in call to
	cp_parser_omp_var_list for to/from clauses.
	* semantics.c (handle_omp_array_sections_1): Extend allowed range
	of base-pointer expressions involving INDIRECT/MEM/ARRAY_REF and
	POINTER_PLUS_EXPR.
	(handle_omp_array_sections): Adjust pointer map generation of
	references.
	(finish_omp_clauses): Extend allowed ranged of expressions
	involving INDIRECT/MEM/ARRAY_REF and POINTER_PLUS_EXPR.

gcc/fortran/ChangeLog:

	* trans-openmp.c (gfc_trans_omp_array_section): Do not generate
	GOMP_MAP_ALWAYS_POINTER map for main array maps of ARRAY_TYPE type.

gcc/ChangeLog:

	* gimplify.c (extract_base_bit_offset): Add 'tree *offsetp' parameter,
	accomodate case where 'offset' return of get_inner_reference is
	non-NULL.
	(is_or_contains_p): Further robustify conditions.
	(omp_target_reorder_clauses): In alloc/to/from sorting phase, also
	move following GOMP_MAP_ALWAYS_POINTER maps along.  Add new sorting
	phase where we make sure pointers with an attach/detach map are ordered
	correctly.
	(gimplify_scan_omp_clauses): Add modifications to avoid creating
	GOMP_MAP_STRUCT and associated alloc map for attach/detach maps.

gcc/testsuite/ChangeLog:

	* c-c++-common/goacc/deep-copy-arrayofstruct.c: Adjust testcase.
	* c-c++-common/gomp/target-enter-data-1.c: New testcase.
	* c-c++-common/gomp/target-implicit-map-2.c: New testcase.

libgomp/ChangeLog:

	* target.c (gomp_map_vars_existing): Make sure attached pointer is
	not overwritten during cross-host/device copying.
	(gomp_update): Likewise.
	(gomp_exit_data): Likewise.
	* testsuite/libgomp.c++/target-11.C: Adjust testcase.
	* testsuite/libgomp.c++/target-12.C: Likewise.
	* testsuite/libgomp.c++/target-15.C: Likewise.
	* testsuite/libgomp.c++/target-16.C: Likewise.
	* testsuite/libgomp.c++/target-17.C: Likewise.
	* testsuite/libgomp.c++/target-21.C: Likewise.
	* testsuite/libgomp.c++/target-23.C: Likewise.
	* testsuite/libgomp.c/target-23.c: Likewise.
	* testsuite/libgomp.c/target-29.c: Likewise.
	* testsuite/libgomp.c-c++-common/target-implicit-map-2.c: New testcase.

10 files changed, 102 insertions, 54 deletions

diff --git a/libgomp/testsuite/libgomp.c++/target-11.C b/libgomp/testsuite/libgomp.c++/target-11.C
index fe99603..87c2980 100644
--- a/libgomp/testsuite/libgomp.c++/target-11.C
+++ b/libgomp/testsuite/libgomp.c++/target-11.C
@@ -23,9 +23,11 @@ foo ()
   e = c + 18;
   D s = { a, b + 2, { 0, a + 16, 0, d }, c + 3, e };
   int err = 0;
-  #pragma omp target map (to:s.v.b[0:z + 7], s.template u[z + 1:z + 4]) \
-		     map (tofrom:s.s[3:3], s. template v. template d[z + 1:z + 3]) \
-		     map (from: s.w[z:4], s.x[1:3], err) private (i)
+  #pragma omp target map (to: s.v.b, s.v.b[0:z + 7])			\
+		     map (s.template u, s.template u[z + 1:z + 4])	\
+		     map (tofrom: s.s, s.s[3:3])			\
+		     map (tofrom: s. template v. template d[z + 1:z + 3])\
+		     map (from: s.w, s.w[z:4], s.x, s.x[1:3], err) private (i)
   {
     err = 0;
     for (i = 0; i < 7; i++)
@@ -80,9 +82,9 @@ main ()
   e = c + 18;
   S s = { a, b + 2, { 0, a + 16, 0, d }, c + 3, e };
   int err = 0;
-  #pragma omp target map (to:s.v.b[0:z + 7], s.u[z + 1:z + 4]) \
-		     map (tofrom:s.s[3:3], s.v.d[z + 1:z + 3]) \
-		     map (from: s.w[z:4], s.x[1:3], err) private (i)
+  #pragma omp target map (to: s.v.b, s.v.b[0:z + 7], s.u, s.u[z + 1:z + 4]) \
+		     map (tofrom: s.s, s.s[3:3], s.v.d[z + 1:z + 3])		\
+		     map (from: s.w, s.w[z:4], s.x, s.x[1:3], err) private (i)
   {
     err = 0;
     for (i = 0; i < 7; i++)
diff --git a/libgomp/testsuite/libgomp.c++/target-12.C b/libgomp/testsuite/libgomp.c++/target-12.C
index 3b4ed57..480e479 100644
--- a/libgomp/testsuite/libgomp.c++/target-12.C
+++ b/libgomp/testsuite/libgomp.c++/target-12.C
@@ -53,7 +53,7 @@ main ()
   int u[10] = { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9 }, err = 0;
   S s = { 9, u + 3, { 10, 11, 12, 13, 14 } };
   int *v = u + 4;
-  #pragma omp target enter data map (to: s.s, s.u[0:5]) map (alloc: s.v[1:3])
+  #pragma omp target enter data map (to: s.s, s.u, s.u[0:5]) map (alloc: s.v[1:3])
   s.s++;
   u[3]++;
   s.v[1]++;
diff --git a/libgomp/testsuite/libgomp.c++/target-15.C b/libgomp/testsuite/libgomp.c++/target-15.C
index 4b320c3..53626b2 100644
--- a/libgomp/testsuite/libgomp.c++/target-15.C
+++ b/libgomp/testsuite/libgomp.c++/target-15.C
@@ -14,7 +14,7 @@ foo (S s)
     d = id;
 
   int err;
-  #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(to: sep) map(from: err)
+  #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(to: sep) map(from: err)
   {
     err = s.a != 11 || s.b[0] != 12 || s.b[1] != 13;
     err |= s.c[1] != 15 || s.c[2] != 16 || s.d[-2] != 18 || s.d[-1] != 19 || s.d[0] != 20;
@@ -48,7 +48,7 @@ foo (S s)
 	  || omp_target_is_present (&s.h, d)
 	  || omp_target_is_present (&s.h[2], d)))
     abort ();
-  #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
   {
     if (!omp_target_is_present (&s.a, d)
 	|| !omp_target_is_present (s.b, d)
@@ -61,8 +61,8 @@ foo (S s)
 	|| !omp_target_is_present (&s.h, d)
 	|| !omp_target_is_present (&s.h[2], d))
       abort ();
-    #pragma omp target update to(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
-    #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
+    #pragma omp target update to(s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
+    #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(from: err)
     {
       err = s.a != 50 || s.b[0] != 49 || s.b[1] != 48;
       err |= s.c[1] != 47 || s.c[2] != 46 || s.d[-2] != 45 || s.d[-1] != 44 || s.d[0] != 43;
@@ -73,7 +73,7 @@ foo (S s)
       s.e = 25; s.f[0] = 26; s.f[1] = 27; s.g[1] = 28; s.g[2] = 29;
       s.h[2] = 30; s.h[3] = 31; s.h[4] = 32;
     }
-    #pragma omp target update from(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+    #pragma omp target update from(s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
   }
   if (sep
       && (omp_target_is_present (&s.a, d)
@@ -97,7 +97,7 @@ foo (S s)
   s.c[1] = 36; s.c[2] = 37; s.d[-2] = 38; s.d[-1] = 39; s.d[0] = 40;
   s.e = 41; s.f[0] = 42; s.f[1] = 43; s.g[1] = 44; s.g[2] = 45;
   s.h[2] = 46; s.h[3] = 47; s.h[4] = 48;
-  #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
   if (!omp_target_is_present (&s.a, d)
       || !omp_target_is_present (s.b, d)
       || !omp_target_is_present (&s.c[1], d)
@@ -109,8 +109,8 @@ foo (S s)
       || !omp_target_is_present (&s.h, d)
       || !omp_target_is_present (&s.h[2], d))
     abort ();
-  #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
-  #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
+  #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
+  #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(from: err)
   {
     err = s.a != 33 || s.b[0] != 34 || s.b[1] != 35;
     err |= s.c[1] != 36 || s.c[2] != 37 || s.d[-2] != 38 || s.d[-1] != 39 || s.d[0] != 40;
@@ -121,7 +121,7 @@ foo (S s)
     s.e = 31; s.f[0] = 40; s.f[1] = 39; s.g[1] = 38; s.g[2] = 37;
     s.h[2] = 36; s.h[3] = 35; s.h[4] = 34;
   }
-  #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
   if (!omp_target_is_present (&s.a, d)
       || !omp_target_is_present (s.b, d)
       || !omp_target_is_present (&s.c[1], d)
@@ -133,7 +133,7 @@ foo (S s)
       || !omp_target_is_present (&s.h, d)
       || !omp_target_is_present (&s.h[2], d))
     abort ();
-  #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
   if (sep
       && (omp_target_is_present (&s.a, d)
 	  || omp_target_is_present (s.b, d)
diff --git a/libgomp/testsuite/libgomp.c++/target-16.C b/libgomp/testsuite/libgomp.c++/target-16.C
index cd102d9..b8be7cc 100644
--- a/libgomp/testsuite/libgomp.c++/target-16.C
+++ b/libgomp/testsuite/libgomp.c++/target-16.C
@@ -16,7 +16,7 @@ foo (S<C, I, L, UC, SH> s)
     d = id;
 
   int err;
-  #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(to: sep) map(from: err)
+  #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(to: sep) map(from: err)
   {
     err = s.a != 11 || s.b[0] != 12 || s.b[1] != 13;
     err |= s.c[1] != 15 || s.c[2] != 16 || s.d[-2] != 18 || s.d[-1] != 19 || s.d[0] != 20;
@@ -50,7 +50,7 @@ foo (S<C, I, L, UC, SH> s)
 	  || omp_target_is_present (&s.h, d)
 	  || omp_target_is_present (&s.h[2], d)))
     abort ();
-  #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
   {
     if (!omp_target_is_present (&s.a, d)
 	|| !omp_target_is_present (s.b, d)
@@ -63,8 +63,8 @@ foo (S<C, I, L, UC, SH> s)
 	|| !omp_target_is_present (&s.h, d)
 	|| !omp_target_is_present (&s.h[2], d))
       abort ();
-    #pragma omp target update to(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
-    #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
+    #pragma omp target update to(s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
+    #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(from: err)
     {
       err = s.a != 50 || s.b[0] != 49 || s.b[1] != 48;
       err |= s.c[1] != 47 || s.c[2] != 46 || s.d[-2] != 45 || s.d[-1] != 44 || s.d[0] != 43;
@@ -75,7 +75,7 @@ foo (S<C, I, L, UC, SH> s)
       s.e = 25; s.f[0] = 26; s.f[1] = 27; s.g[1] = 28; s.g[2] = 29;
       s.h[2] = 30; s.h[3] = 31; s.h[4] = 32;
     }
-    #pragma omp target update from(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+    #pragma omp target update from(s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
   }
   if (sep
       && (omp_target_is_present (&s.a, d)
@@ -99,7 +99,7 @@ foo (S<C, I, L, UC, SH> s)
   s.c[1] = 36; s.c[2] = 37; s.d[-2] = 38; s.d[-1] = 39; s.d[0] = 40;
   s.e = 41; s.f[0] = 42; s.f[1] = 43; s.g[1] = 44; s.g[2] = 45;
   s.h[2] = 46; s.h[3] = 47; s.h[4] = 48;
-  #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
   if (!omp_target_is_present (&s.a, d)
       || !omp_target_is_present (s.b, d)
       || !omp_target_is_present (&s.c[1], d)
@@ -111,8 +111,8 @@ foo (S<C, I, L, UC, SH> s)
       || !omp_target_is_present (&s.h, d)
       || !omp_target_is_present (&s.h[2], d))
     abort ();
-  #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
-  #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
+  #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
+  #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(from: err)
   {
     err = s.a != 33 || s.b[0] != 34 || s.b[1] != 35;
     err |= s.c[1] != 36 || s.c[2] != 37 || s.d[-2] != 38 || s.d[-1] != 39 || s.d[0] != 40;
@@ -123,7 +123,7 @@ foo (S<C, I, L, UC, SH> s)
     s.e = 31; s.f[0] = 40; s.f[1] = 39; s.g[1] = 38; s.g[2] = 37;
     s.h[2] = 36; s.h[3] = 35; s.h[4] = 34;
   }
-  #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
   if (!omp_target_is_present (&s.a, d)
       || !omp_target_is_present (s.b, d)
       || !omp_target_is_present (&s.c[1], d)
@@ -135,7 +135,7 @@ foo (S<C, I, L, UC, SH> s)
       || !omp_target_is_present (&s.h, d)
       || !omp_target_is_present (&s.h[2], d))
     abort ();
-  #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
   if (sep
       && (omp_target_is_present (&s.a, d)
 	  || omp_target_is_present (s.b, d)
diff --git a/libgomp/testsuite/libgomp.c++/target-17.C b/libgomp/testsuite/libgomp.c++/target-17.C
index d81ff19..f97476a 100644
--- a/libgomp/testsuite/libgomp.c++/target-17.C
+++ b/libgomp/testsuite/libgomp.c++/target-17.C
@@ -16,7 +16,7 @@ foo (S<C, I, L, UCR, CAR, SH, IPR> s)
     d = id;
 
   int err;
-  #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(to: sep) map(from: err)
+  #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(to: sep) map(from: err)
   {
     err = s.a != 11 || s.b[0] != 12 || s.b[1] != 13;
     err |= s.c[1] != 15 || s.c[2] != 16 || s.d[-2] != 18 || s.d[-1] != 19 || s.d[0] != 20;
@@ -50,7 +50,7 @@ foo (S<C, I, L, UCR, CAR, SH, IPR> s)
 	  || omp_target_is_present (&s.h, d)
 	  || omp_target_is_present (&s.h[2], d)))
     abort ();
-  #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
   {
     if (!omp_target_is_present (&s.a, d)
 	|| !omp_target_is_present (s.b, d)
@@ -63,8 +63,8 @@ foo (S<C, I, L, UCR, CAR, SH, IPR> s)
 	|| !omp_target_is_present (&s.h, d)
 	|| !omp_target_is_present (&s.h[2], d))
       abort ();
-    #pragma omp target update to(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
-    #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
+    #pragma omp target update to(s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
+    #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(from: err)
     {
       err = s.a != 50 || s.b[0] != 49 || s.b[1] != 48;
       err |= s.c[1] != 47 || s.c[2] != 46 || s.d[-2] != 45 || s.d[-1] != 44 || s.d[0] != 43;
@@ -75,7 +75,7 @@ foo (S<C, I, L, UCR, CAR, SH, IPR> s)
       s.e = 25; s.f[0] = 26; s.f[1] = 27; s.g[1] = 28; s.g[2] = 29;
       s.h[2] = 30; s.h[3] = 31; s.h[4] = 32;
     }
-    #pragma omp target update from(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+    #pragma omp target update from(s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
   }
   if (sep
       && (omp_target_is_present (&s.a, d)
@@ -99,7 +99,7 @@ foo (S<C, I, L, UCR, CAR, SH, IPR> s)
   s.c[1] = 36; s.c[2] = 37; s.d[-2] = 38; s.d[-1] = 39; s.d[0] = 40;
   s.e = 41; s.f[0] = 42; s.f[1] = 43; s.g[1] = 44; s.g[2] = 45;
   s.h[2] = 46; s.h[3] = 47; s.h[4] = 48;
-  #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
   if (!omp_target_is_present (&s.a, d)
       || !omp_target_is_present (s.b, d)
       || !omp_target_is_present (&s.c[1], d)
@@ -111,8 +111,8 @@ foo (S<C, I, L, UCR, CAR, SH, IPR> s)
       || !omp_target_is_present (&s.h, d)
       || !omp_target_is_present (&s.h[2], d))
     abort ();
-  #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
-  #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
+  #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
+  #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(from: err)
   {
     err = s.a != 33 || s.b[0] != 34 || s.b[1] != 35;
     err |= s.c[1] != 36 || s.c[2] != 37 || s.d[-2] != 38 || s.d[-1] != 39 || s.d[0] != 40;
@@ -123,7 +123,7 @@ foo (S<C, I, L, UCR, CAR, SH, IPR> s)
     s.e = 31; s.f[0] = 40; s.f[1] = 39; s.g[1] = 38; s.g[2] = 37;
     s.h[2] = 36; s.h[3] = 35; s.h[4] = 34;
   }
-  #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
   if (!omp_target_is_present (&s.a, d)
       || !omp_target_is_present (s.b, d)
       || !omp_target_is_present (&s.c[1], d)
@@ -135,7 +135,7 @@ foo (S<C, I, L, UCR, CAR, SH, IPR> s)
       || !omp_target_is_present (&s.h, d)
       || !omp_target_is_present (&s.h[2], d))
     abort ();
-  #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+  #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
   if (sep
       && (omp_target_is_present (&s.a, d)
 	  || omp_target_is_present (s.b, d)
diff --git a/libgomp/testsuite/libgomp.c++/target-21.C b/libgomp/testsuite/libgomp.c++/target-21.C
index 21a2f29..da17b57 100644
--- a/libgomp/testsuite/libgomp.c++/target-21.C
+++ b/libgomp/testsuite/libgomp.c++/target-21.C
@@ -7,7 +7,7 @@ void
 foo (S s)
 {
   int err;
-  #pragma omp target map (s.x[0:N], s.y[0:N]) map (s.t.t[16:3]) map (from: err)
+  #pragma omp target map (s.x[0:N], s.y, s.y[0:N]) map (s.t.t[16:3]) map (from: err)
   {
     err = s.x[2] != 28 || s.y[2] != 37 || s.t.t[17] != 81;
     s.x[2]++;
@@ -38,7 +38,7 @@ void
 foo2 (S &s)
 {
   int err;
-  #pragma omp target map (s.x[N:10], s.y[N:10]) map (from: err) map (s.t.t[N+16:N+3])
+  #pragma omp target map (s.x[N:10], s.y, s.y[N:10]) map (from: err) map (s.t.t[N+16:N+3])
   {
     err = s.x[2] != 30 || s.y[2] != 38 || s.t.t[17] != 81;
     s.x[2]++;
@@ -69,7 +69,7 @@ void
 foo3 (U s)
 {
   int err;
-  #pragma omp target map (s.x[0:10], s.y[0:10]) map (from: err) map (s.t.t[16:3])
+  #pragma omp target map (s.x[0:10], s.y, s.y[0:10]) map (from: err) map (s.t.t[16:3])
   {
     err = s.x[2] != 32 || s.y[2] != 39 || s.t.t[17] != 82;
     s.x[2]++;
@@ -100,7 +100,7 @@ void
 foo4 (U &s)
 {
   int err;
-  #pragma omp target map (s.x[0:10], s.y[0:10]) map (from: err) map (s.t.t[16:3])
+  #pragma omp target map (s.x[0:10], s.y, s.y[0:10]) map (from: err) map (s.t.t[16:3])
   {
     err = s.x[2] != 34 || s.y[2] != 40 || s.t.t[17] != 82;
     s.x[2]++;
diff --git a/libgomp/testsuite/libgomp.c++/target-23.C b/libgomp/testsuite/libgomp.c++/target-23.C
index d4f9ff3..63d3436 100644
--- a/libgomp/testsuite/libgomp.c++/target-23.C
+++ b/libgomp/testsuite/libgomp.c++/target-23.C
@@ -16,13 +16,13 @@ main (void)
     s->data[i] = 0;
 
   #pragma omp target enter data map(to: s)
-  #pragma omp target enter data map(to: s->data[:SZ])
+  #pragma omp target enter data map(to: s->data, s->data[:SZ])
   #pragma omp target
   {
     for (int i = 0; i < SZ; i++)
       s->data[i] = i;
   }
-  #pragma omp target exit data map(from: s->data[:SZ])
+  #pragma omp target exit data map(from: s->data, s->data[:SZ])
   #pragma omp target exit data map(from: s)
 
   for (int i = 0; i < SZ; i++)
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-2.c b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-2.c
new file mode 100644
index 0000000..974a978
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-2.c
@@ -0,0 +1,46 @@
+#include <stdlib.h>
+
+#define N 10
+
+struct S
+{
+  int a, b;
+  int *ptr;
+  int c, d;
+};
+
+int
+main (void)
+{
+  struct S a;
+  a.ptr = (int *) malloc (sizeof (int) * N);
+
+  for (int i = 0; i < N; i++)
+    a.ptr[i] = 0;
+
+  #pragma omp target enter data map(to: a.ptr, a.ptr[:N])
+
+  #pragma omp target
+  for (int i = 0; i < N; i++)
+    a.ptr[i] += 1;
+
+  #pragma omp target update from(a.ptr[:N])
+
+  for (int i = 0; i < N; i++)
+    if (a.ptr[i] != 1)
+      abort ();
+
+  #pragma omp target map(a.ptr[:N])
+  for (int i = 0; i < N; i++)
+    a.ptr[i] += 1;
+
+  #pragma omp target update from(a.ptr[:N])
+
+  for (int i = 0; i < N; i++)
+    if (a.ptr[i] != 2)
+      abort ();
+
+  #pragma omp target exit data map(from:a.ptr, a.ptr[:N])
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/target-23.c b/libgomp/testsuite/libgomp.c/target-23.c
index fb1532a..d56b13a 100644
--- a/libgomp/testsuite/libgomp.c/target-23.c
+++ b/libgomp/testsuite/libgomp.c/target-23.c
@@ -8,7 +8,7 @@ main ()
   int u[10] = { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9 }, err = 0;
   struct S s = { 9, u + 3, { 10, 11, 12, 13, 14 } };
   int *v = u + 4;
-  #pragma omp target enter data map (to: s.s, s.u[0:5]) map (alloc: s.v[1:3])
+  #pragma omp target enter data map (to: s.s, s.u, s.u[0:5]) map (alloc: s.v[1:3])
   s.s++;
   u[3]++;
   s.v[1]++;
diff --git a/libgomp/testsuite/libgomp.c/target-29.c b/libgomp/testsuite/libgomp.c/target-29.c
index e5095a1..4a28664 100644
--- a/libgomp/testsuite/libgomp.c/target-29.c
+++ b/libgomp/testsuite/libgomp.c/target-29.c
@@ -14,7 +14,7 @@ foo (struct S s)
     d = id;
 
   int err;
-  #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d[-2:3]) map(to: sep) map(from: err)
+  #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d, s.d[-2:3]) map(to: sep) map(from: err)
   {
     err = s.a != 11 || s.b[0] != 12 || s.b[1] != 13;
     err |= s.c[1] != 15 || s.c[2] != 16 || s.d[-2] != 18 || s.d[-1] != 19 || s.d[0] != 20;
@@ -35,7 +35,7 @@ foo (struct S s)
 	  || omp_target_is_present (s.d, d)
 	  || omp_target_is_present (&s.d[-2], d)))
     abort ();
-  #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3])
+  #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
   {
     if (!omp_target_is_present (&s.a, d)
 	|| !omp_target_is_present (s.b, d)
@@ -43,15 +43,15 @@ foo (struct S s)
 	|| !omp_target_is_present (s.d, d)
 	|| !omp_target_is_present (&s.d[-2], d))
       abort ();
-    #pragma omp target update to(s.a, s.b, s.c[1:2], s.d[-2:3])
-    #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3]) map(from: err)
+    #pragma omp target update to(s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
+    #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3]) map(from: err)
     {
       err = s.a != 50 || s.b[0] != 49 || s.b[1] != 48;
       err |= s.c[1] != 47 || s.c[2] != 46 || s.d[-2] != 45 || s.d[-1] != 44 || s.d[0] != 43;
       s.a = 17; s.b[0] = 18; s.b[1] = 19;
       s.c[1] = 20; s.c[2] = 21; s.d[-2] = 22; s.d[-1] = 23; s.d[0] = 24;
     }
-    #pragma omp target update from(s.a, s.b, s.c[1:2], s.d[-2:3])
+    #pragma omp target update from(s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
   }
   if (sep
       && (omp_target_is_present (&s.a, d)
@@ -66,29 +66,29 @@ foo (struct S s)
   if (err) abort ();
   s.a = 33; s.b[0] = 34; s.b[1] = 35;
   s.c[1] = 36; s.c[2] = 37; s.d[-2] = 38; s.d[-1] = 39; s.d[0] = 40;
-  #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3])
+  #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
   if (!omp_target_is_present (&s.a, d)
       || !omp_target_is_present (s.b, d)
       || !omp_target_is_present (&s.c[1], d)
       || !omp_target_is_present (s.d, d)
       || !omp_target_is_present (&s.d[-2], d))
     abort ();
-  #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d[-2:3])
-  #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3]) map(from: err)
+  #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
+  #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3]) map(from: err)
   {
     err = s.a != 33 || s.b[0] != 34 || s.b[1] != 35;
     err |= s.c[1] != 36 || s.c[2] != 37 || s.d[-2] != 38 || s.d[-1] != 39 || s.d[0] != 40;
     s.a = 49; s.b[0] = 48; s.b[1] = 47;
     s.c[1] = 46; s.c[2] = 45; s.d[-2] = 44; s.d[-1] = 43; s.d[0] = 42;
   }
-  #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d[-2:3])
+  #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
   if (!omp_target_is_present (&s.a, d)
       || !omp_target_is_present (s.b, d)
       || !omp_target_is_present (&s.c[1], d)
       || !omp_target_is_present (s.d, d)
       || !omp_target_is_present (&s.d[-2], d))
     abort ();
-  #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d[-2:3])
+  #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
   if (sep
       && (omp_target_is_present (&s.a, d)
 	  || omp_target_is_present (s.b, d)