https://gcc.gnu.org/g:87548cbb46416486888daad14c606dffa8e58204
commit 87548cbb46416486888daad14c606dffa8e58204
Author: Julian Brown <jul...@codesourcery.com>
Date: Sun Apr 20 00:00:54 2025 +0000
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 <tschwi...@baylibre.com>
Co-Authored-By: Sandra Loosemore <sloosem...@baylibre.com>
Diff:
---
gcc/gimplify.cc | 5 +----
.../c-c++-common/goacc/combined-reduction.c | 2 +-
gcc/testsuite/c-c++-common/goacc/implied-copy-1.c | 4 ++--
gcc/testsuite/c-c++-common/goacc/reduction-1.c | 4 ++--
gcc/testsuite/c-c++-common/goacc/reduction-10.c | 9 ++++----
gcc/testsuite/c-c++-common/goacc/reduction-2.c | 4 ++--
gcc/testsuite/c-c++-common/goacc/reduction-3.c | 4 ++--
gcc/testsuite/c-c++-common/goacc/reduction-4.c | 4 ++--
gcc/testsuite/gfortran.dg/goacc/common-block-3.f90 | 8 +++----
gcc/testsuite/gfortran.dg/goacc/implied-copy-1.f90 | 4 ++--
gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90 | 2 +-
.../goacc/private-explicit-kernels-1.f95 | 8 +++----
.../goacc/private-predetermined-kernels-1.f95 | 8 +++----
include/gomp-constants.h | 3 ++-
.../libgomp.oacc-c-c++-common/implicit-mapping-1.c | 25 ++++++++++++++++++++++
15 files changed, 59 insertions(+), 35 deletions(-)
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 00fb34f71827..a83af5fa4d3e 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 0a541402d9d4..75fef98b0bb9 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 ae06339dc2d8..34ce9b0713db 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 35bfc868708e..7f3e3e33cbda 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 579aa561479d..cd0d58f81402 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 9dba035adb6e..89bc1644dea5 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 669cd4381133..9cca625a925f 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 5c3dfb191721..80b32476062f 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 4861a643ecac..dc222c6feaa6 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 7f07c8afbece..b7f2faacc052 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 150f9304e461..4b2b49f119eb 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 5d563d226b0c..f54b2448f0be 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 12a7854526a9..49b41a55a211 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 217f45747305..1bf5a9111b17 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 000000000000..ed0ab94cd8fe
--- /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;
+}
