OpenACC: Improve implicit mapping for non-lexically nested offload regions

This patch enables use of the OMP_CLAUSE_RUNTIME_IMPLICIT_P flag for
OpenACC.

This allows code like this to work correctly:

  int arr[100];
  [...]
  #pragma acc enter data copyin(arr[20:10])

  /* No explicit mapping of 'arr' here.  */
  #pragma acc parallel
  { /* use of arr[20:10]... */ }

  #pragma acc exit data copyout(arr[20:10])

Otherwise, the implicit "copy" ("present_or_copy") on the parallel
corresponds to the whole array, and that fails at runtime when the
subarray is mapped.

The numbering of the GOMP_MAP_IMPLICIT bit clashes with the OpenACC
"non-contiguous" dynamic array support, so the GOMP_MAP_NONCONTIG_ARRAY_P
macro has been adjusted to account for that.

gcc/
	* gimplify.cc (gimplify_adjust_omp_clauses_1): Set
	OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P for OpenACC also.

gcc/testsuite/
	* c-c++-common/goacc/combined-reduction.c: Adjust scan output.
	* c-c++-common/goacc/implied-copy-1.c: Likewise.
	* c-c++-common/goacc/reduction-1.c: Adjust patterns.
	* c-c++-common/goacc/reduction-2.c: Likewise.
	* c-c++-common/goacc/reduction-3.c: Likewise.
	* c-c++-common/goacc/reduction-4.c: Likewise.
	* c-c++-common/goacc/reduction-10.c: Likewise.
	* gfortran.dg/goacc/common-block-3.f90: Likewise.
	* gfortran.dg/goacc/implied-copy-1.f90: Likewise.
	* gfortran.dg/goacc/loop-tree-1.f90: Likewise.
	* gfortran.dg/goacc/private-explicit-kernels-1.f95: Likewise.
	* gfortran.dg/goacc/private-predetermined-kernels-1.f95: Likewise.

include/
	* gomp-constants.h (GOMP_MAP_NONCONTIG_ARRAY_P): Tweak condition.

libgomp/
	* testsuite/libgomp.oacc-c-c++-common/implicit-mapping-1.c: New test.

Co-Authored-By: Thomas Schwinge <tschwinge@baylibre.com>
Co-Authored-By: Sandra Loosemore <sloosemore@baylibre.com>

15 files changed, 59 insertions, 35 deletions

diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 00fb34f..a83af5f 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -14858,10 +14858,7 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
 	  gcc_unreachable ();
 	}
       OMP_CLAUSE_SET_MAP_KIND (clause, kind);
-      /* Setting of the implicit flag for the runtime is currently disabled for
-	 OpenACC.  */
-      if ((gimplify_omp_ctxp->region_type & ORT_ACC) == 0)
-	OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P (clause) = 1;
+      OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P (clause) = 1;
       if (DECL_SIZE (decl)
 	  && !poly_int_tree_p (DECL_SIZE (decl)))
 	{
diff --git a/gcc/testsuite/c-c++-common/goacc/combined-reduction.c b/gcc/testsuite/c-c++-common/goacc/combined-reduction.c
index 0a54140..75fef98 100644
--- a/gcc/testsuite/c-c++-common/goacc/combined-reduction.c
+++ b/gcc/testsuite/c-c++-common/goacc/combined-reduction.c
@@ -33,7 +33,7 @@ main ()
 
 /* { dg-final { scan-tree-dump-times "omp target oacc_parallel reduction.+:v1. map.tofrom:v1" 1 "gimple" } } */
 /* { dg-final { scan-tree-dump-times "acc loop reduction.+:v1. private.i." 1 "gimple" } } */
-/* { dg-final { scan-tree-dump-times "omp target oacc_kernels map.force_tofrom:n .len: 4.. map.force_tofrom:v1 .len: 4.." 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "omp target oacc_kernels map.force_tofrom:n .len: 4. .runtime_implicit.. map.force_tofrom:v1 .len: 4. .runtime_implicit.." 1 "gimple" } } */
 /* { dg-final { scan-tree-dump-times "acc loop reduction.+:v1. private.i." 1 "gimple" } } */
 /* { dg-final { scan-tree-dump-times "omp target oacc_serial reduction.+:v1. map.tofrom:v1" 1 "gimple" } } */
 /* { dg-final { scan-tree-dump-times "acc loop reduction.+:v1. private.i." 1 "gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/implied-copy-1.c b/gcc/testsuite/c-c++-common/goacc/implied-copy-1.c
index ae06339..34ce9b0 100644
--- a/gcc/testsuite/c-c++-common/goacc/implied-copy-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/implied-copy-1.c
@@ -27,7 +27,7 @@ void test1 (void)
   }
 }
 
-/* { dg-final { scan-tree-dump-times "map\\(force_tofrom:sum \\\[len: \[0-9\]+\\\]\\)" 1 "gimple" } } */
-/* { dg-final { scan-tree-dump-times "map\\(force_tofrom:prod \\\[len: \[0-9\]+\\\]\\)" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(force_tofrom:sum \\\[len: \[0-9\]+\\\].*\\)" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(force_tofrom:prod \\\[len: \[0-9\]+\\\].*\\)" 1 "gimple" } } */
 /* { dg-final { scan-tree-dump-times "map\\(tofrom:sum \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */
 /* { dg-final { scan-tree-dump-times "map\\(tofrom:prod \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-1.c b/gcc/testsuite/c-c++-common/goacc/reduction-1.c
index 35bfc86..7f3e3e3 100644
--- a/gcc/testsuite/c-c++-common/goacc/reduction-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/reduction-1.c
@@ -68,5 +68,5 @@ main(void)
 }
 
 /* Check that default copy maps are generated for loop reductions.  */
-/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: \[0-9\]+\\\]\\)" 7 "gimple" } } */
-/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: \[0-9\]+\\\] \\\[runtime_implicit\\\]\\)" 7 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: \[0-9\]+\\\] \\\[runtime_implicit\\\]\\)" 2 "gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-10.c b/gcc/testsuite/c-c++-common/goacc/reduction-10.c
index 579aa56..cd0d58f 100644
--- a/gcc/testsuite/c-c++-common/goacc/reduction-10.c
+++ b/gcc/testsuite/c-c++-common/goacc/reduction-10.c
@@ -87,7 +87,8 @@ main(void)
 
 /* Check that default copy maps are generated for loop reductions.  */
 /* { dg-final { scan-tree-dump-times "reduction..:result. map.tofrom:result .len: 4.." 1 "gimple" } } */
-/* { dg-final { scan-tree-dump-times "oacc_parallel map.tofrom:result .len: 4.." 2 "gimple" } } */
-/* { dg-final { scan-tree-dump-times "map.tofrom:array .len: 4000.. firstprivate.result." 3 "gimple" } } */
-/* { dg-final { scan-tree-dump-times "map.tofrom:result .len: 4.. map.tofrom:array .len: 4000.." 1 "gimple" } } */
-/* { dg-final { scan-tree-dump-times "map.tofrom:array .len: 4000.. map.force_tofrom:result .len: 4.." 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times {oacc_parallel map\(tofrom:result \[len: 4\]\)} 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "oacc_parallel map.tofrom:result .len: 4. .runtime_implicit.." 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map.tofrom:array .len: 4000. .runtime_implicit.. firstprivate.result." 3 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map.tofrom:result .len: 4.. map.tofrom:array .len: 4000. .runtime_implicit.." 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map.tofrom:array .len: 4000. .runtime_implicit.. map.force_tofrom:result .len: 4. .runtime_implicit.." 1 "gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-2.c b/gcc/testsuite/c-c++-common/goacc/reduction-2.c
index 9dba035..89bc164 100644
--- a/gcc/testsuite/c-c++-common/goacc/reduction-2.c
+++ b/gcc/testsuite/c-c++-common/goacc/reduction-2.c
@@ -50,5 +50,5 @@ main(void)
 }
 
 /* Check that default copy maps are generated for loop reductions.  */
-/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: \[0-9\]+\\\]\\)" 4 "gimple" } } */
-/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: \[0-9\]+\\\] \\\[runtime_implicit\\\]\\)" 4 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: \[0-9\]+\\\] \\\[runtime_implicit\\\]\\)" 2 "gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-3.c b/gcc/testsuite/c-c++-common/goacc/reduction-3.c
index 669cd43..9cca625 100644
--- a/gcc/testsuite/c-c++-common/goacc/reduction-3.c
+++ b/gcc/testsuite/c-c++-common/goacc/reduction-3.c
@@ -50,5 +50,5 @@ main(void)
 }
 
 /* Check that default copy maps are generated for loop reductions.  */
-/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: \[0-9\]+\\\]\\)" 4 "gimple" } } */
-/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: \[0-9\]+\\\] \\\[runtime_implicit\\\]\\)" 4 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: \[0-9\]+\\\] \\\[runtime_implicit\\\]\\)" 2 "gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-4.c b/gcc/testsuite/c-c++-common/goacc/reduction-4.c
index 5c3dfb1..80b3247 100644
--- a/gcc/testsuite/c-c++-common/goacc/reduction-4.c
+++ b/gcc/testsuite/c-c++-common/goacc/reduction-4.c
@@ -38,5 +38,5 @@ main(void)
 }
 
 /* Check that default copy maps are generated for loop reductions.  */
-/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */
-/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: \[0-9\]+\\\] \\\[runtime_implicit\\\]\\)" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: \[0-9\]+\\\] \\\[runtime_implicit\\\]\\)" 2 "gimple" } } */
diff --git a/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90 b/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90
index 4861a64..dc222c6 100644
--- a/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90
@@ -52,13 +52,13 @@ end program main
 ! { dg-final { scan-tree-dump-times "omp target oacc_parallel .*map\\(tofrom:c \\\[len: 4\\\]\\)" 1 "omplower" } }
 
 ! { dg-final { scan-tree-dump-times "omp target oacc_data_kernels .*map\\(force_tofrom:i \\\[len: 4\\\]\\)" 1 "omplower" } }
-!   { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_present:i \\\[len: 4\\\]\\)" 1 "omplower" } }
+!   { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_present:i \\\[len: 4\\\].*\\)" 1 "omplower" } }
 ! { dg-final { scan-tree-dump-times "omp target oacc_data_kernels .*map\\(tofrom:x \\\[len: 400\\\]\\)" 1 "omplower" } }
-!   { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_present:x \\\[len: 400\\\]\\)" 1 "omplower" } }
+!   { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_present:x \\\[len: 400\\\].*\\)" 1 "omplower" } }
 ! { dg-final { scan-tree-dump-times "omp target oacc_data_kernels .*map\\(tofrom:y \\\[len: 400\\\]\\\)" 1 "omplower" } }
-!   { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_present:y \\\[len: 400\\\]\\\)" 1 "omplower" } }
+!   { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_present:y \\\[len: 400\\\].*\\\)" 1 "omplower" } }
 ! { dg-final { scan-tree-dump-times "omp target oacc_data_kernels .*map\\(force_tofrom:c \\\[len: 4\\\]\\)" 1 "omplower" } }
-!   { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_present:c \\\[len: 4\\\]\\)" 1 "omplower" } }
+!   { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_present:c \\\[len: 4\\\].*\\)" 1 "omplower" } }
 
 ! { dg-final { scan-tree-dump-times "omp target oacc_serial .*map\\(tofrom:a \\\[len: 400\\\]\\)" 1 "omplower" } }
 ! { dg-final { scan-tree-dump-times "omp target oacc_serial .*map\\(tofrom:b \\\[len: 400\\\]\\\)" 1 "omplower" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/implied-copy-1.f90 b/gcc/testsuite/gfortran.dg/goacc/implied-copy-1.f90
index 7f07c8a..b7f2faa 100644
--- a/gcc/testsuite/gfortran.dg/goacc/implied-copy-1.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/implied-copy-1.f90
@@ -29,7 +29,7 @@ subroutine test
   !$acc end kernels loop
 end subroutine test
 
-! { dg-final { scan-tree-dump-times "map\\(force_tofrom:s \\\[len: \[0-9\]+\\\]\\)" 1 "gimple" } } 
-! { dg-final { scan-tree-dump-times "map\\(force_tofrom:p \\\[len: \[0-9\]+\\\]\\)" 1 "gimple" } } 
+! { dg-final { scan-tree-dump-times "map\\(force_tofrom:s \\\[len: \[0-9\]+\\\].*\\)" 1 "gimple" } } 
+! { dg-final { scan-tree-dump-times "map\\(force_tofrom:p \\\[len: \[0-9\]+\\\].*\\)" 1 "gimple" } } 
 ! { dg-final { scan-tree-dump-times "map\\(tofrom:s \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } 
 ! { dg-final { scan-tree-dump-times "map\\(tofrom:p \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } 
diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90 b/gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90
index 150f930..4b2b49f 100644
--- a/gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90
@@ -44,4 +44,4 @@ end program test
 
 ! { dg-final { scan-tree-dump-times "private\\(m\\)" 1 "original" } } 
 ! { dg-final { scan-tree-dump-times "reduction\\(\\+:sum\\)" 1 "original" } } 
-! { dg-final { scan-tree-dump-times "map\\(tofrom:sum \\\[len: \[0-9\]+\\\]\\)" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(tofrom:sum \\\[len: \[0-9\]+\\\] \\\[runtime_implicit\\\]\\)" 1 "gimple" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/private-explicit-kernels-1.f95 b/gcc/testsuite/gfortran.dg/goacc/private-explicit-kernels-1.f95
index 5d563d2..f54b244 100644
--- a/gcc/testsuite/gfortran.dg/goacc/private-explicit-kernels-1.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/private-explicit-kernels-1.f95
@@ -30,7 +30,7 @@ program test
   !$acc kernels ! Explicit "private(i0_1)" clause cannot be specified here.
   ! { dg-final { scan-tree-dump-times "private\\(i0_1\\)" 1 "original" { xfail *-*-* } } } ! PR90067
   ! { dg-final { scan-tree-dump-times "private\\(i0_1\\)" 1 "gimple" { xfail *-*-* } } } ! PR90067
-  ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i0_1 \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+  ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i0_1 \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { target *-*-* } } }
   do i0_1 = 1, 100
   end do
   !$acc end kernels
@@ -40,7 +40,7 @@ program test
   ! { dg-final { scan-tree-dump-times "private\\(j0_2\\)" 1 "original" { xfail *-*-* } } } ! PR90067
   ! { dg-final { scan-tree-dump-times "private\\(i0_2\\)" 1 "gimple" { xfail *-*-* } } } ! PR90067
   ! { dg-final { scan-tree-dump-times "private\\(j0_2\\)" 1 "gimple" { xfail *-*-* } } } ! PR90067
-  ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:j0_2 \\\[len: \[0-9\]+\\\]\\) map\\(force_tofrom:i0_2 \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+  ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:j0_2 \\\[len: \[0-9\]+\\\]\\) map\\(force_tofrom:i0_2 \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { target *-*-* } } }
   do i0_2 = 1, 100
      do j0_2 = 1, 100
      end do
@@ -82,7 +82,7 @@ program test
   !$acc kernels ! Explicit "private(i2_2_s)" clause cannot be specified here.
   ! { dg-final { scan-tree-dump-times "private\\(i2_2_s\\)" 1 "original" { xfail *-*-* } } } ! PR90067
   ! { dg-final { scan-tree-dump-times "private\\(i2_2_s\\)" 1 "gimple" { xfail *-*-* } } } ! PR90067
-  ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i2_2_s \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+  ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i2_2_s \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { target *-*-* } } }
   do i2_2_s = 1, 100
      !$acc loop private(j2_2_s) independent
      ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_2_s\\) independent" 1 "original" } }
@@ -231,7 +231,7 @@ program test
   !$acc kernels ! Explicit "private(i3_5_s)" clause cannot be specified here.
   ! { dg-final { scan-tree-dump-times "private\\(i3_5_s\\)" 1 "original" { xfail *-*-* } } } ! PR90067
   ! { dg-final { scan-tree-dump-times "private\\(i3_5_s\\)" 1 "gimple" { xfail *-*-* } } } ! PR90067
-  ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i3_5_s \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+  ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i3_5_s \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { target *-*-* } } }
   do i3_5_s = 1, 100
      !$acc loop private(j3_5_s) independent
      ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_5_s\\) independent" 1 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f95 b/gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f95
index 12a7854..49b41a5 100644
--- a/gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f95
@@ -30,7 +30,7 @@ program test
   !$acc kernels
   ! { dg-final { scan-tree-dump-times "private\\(i0_1\\)" 1 "original" { xfail *-*-* } } } ! PR90067
   ! { dg-final { scan-tree-dump-times "private\\(i0_1\\)" 1 "gimple" { xfail *-*-* } } } ! PR90067
-  ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i0_1 \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+  ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i0_1 \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { target *-*-* } } }
   do i0_1 = 1, 100
   end do
   !$acc end kernels
@@ -40,7 +40,7 @@ program test
   ! { dg-final { scan-tree-dump-times "private\\(j0_2\\)" 1 "original" { xfail *-*-* } } } ! PR90067
   ! { dg-final { scan-tree-dump-times "private\\(i0_2\\)" 1 "gimple" { xfail *-*-* } } } ! PR90067
   ! { dg-final { scan-tree-dump-times "private\\(j0_2\\)" 1 "gimple" { xfail *-*-* } } } ! PR90067
-  ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:j0_2 \\\[len: \[0-9\]+\\\]\\) map\\(force_tofrom:i0_2 \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+  ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:j0_2 \\\[len: \[0-9\]+\\\]\\) map\\(force_tofrom:i0_2 \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { target *-*-* } } }
   do i0_2 = 1, 100
      do j0_2 = 1, 100
      end do
@@ -82,7 +82,7 @@ program test
   !$acc kernels
   ! { dg-final { scan-tree-dump-times "private\\(i2_2_s\\)" 1 "original" { xfail *-*-* } } } ! PR90067
   ! { dg-final { scan-tree-dump-times "private\\(i2_2_s\\)" 1 "gimple" { xfail *-*-* } } } ! PR90067
-  ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i2_2_s \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+  ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i2_2_s \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { target *-*-* } } }
   do i2_2_s = 1, 100
      !$acc loop independent
      ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_2_s\\) independent" 1 "original" } }
@@ -231,7 +231,7 @@ program test
   !$acc kernels
   ! { dg-final { scan-tree-dump-times "private\\(i3_5_s\\)" 1 "original" { xfail *-*-* } } } ! PR90067
   ! { dg-final { scan-tree-dump-times "private\\(i3_5_s\\)" 1 "gimple" { xfail *-*-* } } } ! PR90067
-  ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i3_5_s \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+  ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i3_5_s \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { target *-*-* } } }
   do i3_5_s = 1, 100
      !$acc loop independent
      ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_5_s\\) independent" 1 "original" } }
diff --git a/include/gomp-constants.h b/include/gomp-constants.h
index 217f457..1bf5a91 100644
--- a/include/gomp-constants.h
+++ b/include/gomp-constants.h
@@ -276,7 +276,8 @@ enum gomp_map_kind
    || (X) == GOMP_MAP_FORCE_PRESENT)
 
 #define GOMP_MAP_NONCONTIG_ARRAY_P(X) \
-  ((X) & GOMP_MAP_NONCONTIG_ARRAY)
+  (((X) & GOMP_MAP_NONCONTIG_ARRAY) != 0 \
+   && ((X) & GOMP_MAP_FLAG_SPECIAL_4) == 0)
 
 /* Asynchronous behavior.  Keep in sync with
    libgomp/{openacc.h,openacc.f90,openacc_lib.h}:acc_async_t.  */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/implicit-mapping-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/implicit-mapping-1.c
new file mode 100644
index 0000000..ed0ab94
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/implicit-mapping-1.c
@@ -0,0 +1,25 @@
+/* { dg-do run } */
+
+#include <string.h>
+#include <assert.h>
+
+int main(void)
+{
+  int arr[100];
+
+  memset (arr, 0, sizeof (int) * 100);
+
+#pragma acc enter data copyin(arr[30:10])
+
+#pragma acc serial
+/* { dg-warning {using .vector_length \(32\)., ignoring 1} "" { target openacc_nvidia_accel_selected } .-1 } */
+  {
+    arr[33] = 66;
+  }
+
+#pragma acc exit data copyout(arr[30:10])
+
+  assert (arr[33] == 66);
+
+  return 0;
+}