diff mbox series

[PR90048] Fortran OpenACC 'private' clause rejected for predetermined private loop iteration variable (was: [patch,gomp4] make fortran loop variables implicitly private in openacc)

Message ID yxfpwojt9gcd.fsf@hertz.schwinge.homeip.net
State New
Headers show
Series [PR90048] Fortran OpenACC 'private' clause rejected for predetermined private loop iteration variable (was: [patch,gomp4] make fortran loop variables implicitly private in openacc) | expand

Commit Message

Thomas Schwinge April 17, 2019, 8:38 a.m. UTC 
Hi!

On Mon, 11 Aug 2014 16:55:28 -0700, Cesar Philippidis <cesar@codesourcery.com> wrote:
> According to section 2.6.1 in the openacc spec, fortran loop variables
> should be implicitly private like in openmp.

More correctly, they are "predetermined private" (which cannot be
overridden), not "implicit private" (which could be overridden with a
different explicit clause).

> This patch does just so.

But it also introduced PR90048 "Fortran OpenACC 'private' clause rejected
for predetermined private loop iteration variable".

Instead of the patch "don't error on implicitly private induction
variables in gfortran" proposed by Cesar, and then challenged by Jakub, I
have now committed a different patch (more similar to the existing
handling for OpenMP) to trunk in r270406 "[PR90048] Fortran OpenACC
'private' clause rejected for predetermined private loop iteration
variable", see attached.


I have a cleanup patch (for next GCC development stage 1), which will
simply merge the special-case 'gfc_resolve_oacc_blocks' into the generic
'gfc_resolve_omp_parallel_blocks', see attached.


> --- /dev/null
> +++ b/gcc/testsuite/gfortran.dg/goacc/private-1.f95
> @@ -0,0 +1,39 @@
> +! { dg-do compile } 
> +! { dg-additional-options "-fdump-tree-omplower" } 
> +
> +! test for implicit private clauses in do loops
> +
> +program test
> +  implicit none
> +  integer :: i, j, k
> +  logical :: l
> +
> +  !$acc parallel
> +  !$acc loop
> +  do i = 1, 100
> +  end do
> +  !$acc end parallel
> +
> +  !$acc parallel
> +  !$acc loop
> +  do i = 1, 100
> +     do j = 1, 100
> +     end do
> +  end do
> +  !$acc end parallel
> +
> +  !$acc parallel
> +  !$acc loop
> +  do i = 1, 100
> +     do j = 1, 100
> +        do k = 1, 100
> +        end do
> +     end do
> +  end do
> +  !$acc end parallel
> +end program test
> +! { dg-prune-output "unimplemented" }
> +! { dg-final { scan-tree-dump-times "pragma acc parallel" 3 "omplower" } } 
> +! { dg-final { scan-tree-dump-times "private\\(i\\)" 3 "omplower" } } 
> +! { dg-final { scan-tree-dump-times "private\\(j\\)" 2 "omplower" } } 
> +! { dg-final { scan-tree-dump-times "private\\(k\\)" 1 "omplower" } } 

I turned that one and 'gfortran.dg/goacc/private-2.f95' into more
elaborate testcases, committed to trunk in r270405 "[PR90067, PR90114]
Document Fortran OpenACC predetermined private status quo", see attached.


Grüße
 Thomas
diff mbox series

Patch

From 27d00f9196a9ca0d48b0f0934b715307c4b08f5d Mon Sep 17 00:00:00 2001
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Wed, 17 Apr 2019 08:34:10 +0000
Subject: [PATCH] [PR90067, PR90114] Document Fortran OpenACC predetermined
 private status quo

	gcc/testsuite/
	PR fortran/90067
	PR fortran/90114
	* gfortran.dg/goacc/private-1.f95: Remove file.
	* gfortran.dg/goacc/private-2.f95: Likewise.
	* gfortran.dg/goacc/private-predetermined-kernels-1.f95: New file.
	* gfortran.dg/goacc/private-predetermined-parallel-1.f95:
	Likewise.
	* gfortran.dg/goacc/private-predetermined-routine-1.f95: Likewise.

git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@270405 138bc75d-0d04-0410-961f-82ee72b054a4
---
 gcc/testsuite/ChangeLog                       |  11 +
 gcc/testsuite/gfortran.dg/goacc/private-1.f95 |  37 ---
 gcc/testsuite/gfortran.dg/goacc/private-2.f95 |  39 ---
 .../goacc/private-predetermined-kernels-1.f95 | 248 +++++++++++++++++
 .../private-predetermined-parallel-1.f95      | 253 ++++++++++++++++++
 .../goacc/private-predetermined-routine-1.f95 | 142 ++++++++++
 6 files changed, 654 insertions(+), 76 deletions(-)
 delete mode 100644 gcc/testsuite/gfortran.dg/goacc/private-1.f95
 delete mode 100644 gcc/testsuite/gfortran.dg/goacc/private-2.f95
 create mode 100644 gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f95
 create mode 100644 gcc/testsuite/gfortran.dg/goacc/private-predetermined-parallel-1.f95
 create mode 100644 gcc/testsuite/gfortran.dg/goacc/private-predetermined-routine-1.f95

diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog
index 23f77ec781a6..0ed8b3c52837 100644
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,3 +1,14 @@ 
+2019-04-17  Thomas Schwinge  <thomas@codesourcery.com>
+
+	PR fortran/90067
+	PR fortran/90114
+	* gfortran.dg/goacc/private-1.f95: Remove file.
+	* gfortran.dg/goacc/private-2.f95: Likewise.
+	* gfortran.dg/goacc/private-predetermined-kernels-1.f95: New file.
+	* gfortran.dg/goacc/private-predetermined-parallel-1.f95:
+	Likewise.
+	* gfortran.dg/goacc/private-predetermined-routine-1.f95: Likewise.
+
 2019-04-17  Jakub Jelinek  <jakub@redhat.com>
 
 	PR target/89093
diff --git a/gcc/testsuite/gfortran.dg/goacc/private-1.f95 b/gcc/testsuite/gfortran.dg/goacc/private-1.f95
deleted file mode 100644
index 23ce95ad8d24..000000000000
--- a/gcc/testsuite/gfortran.dg/goacc/private-1.f95
+++ /dev/null
@@ -1,37 +0,0 @@ 
-! { dg-do compile }
-! { dg-additional-options "-fdump-tree-omplower" }
-
-! test for implicit private clauses in do loops
-
-program test
-  implicit none
-  integer :: i, j, k
-
-  !$acc parallel
-  !$acc loop
-  do i = 1, 100
-  end do
-  !$acc end parallel
-
-  !$acc parallel
-  !$acc loop
-  do i = 1, 100
-     do j = 1, 100
-     end do
-  end do
-  !$acc end parallel
-
-  !$acc parallel
-  !$acc loop
-  do i = 1, 100
-     do j = 1, 100
-        do k = 1, 100
-        end do
-     end do
-  end do
-  !$acc end parallel
-end program test
-! { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel" 3 "omplower" } }
-! { dg-final { scan-tree-dump-times "private\\(i\\)" 3 "omplower" } }
-! { dg-final { scan-tree-dump-times "private\\(j\\)" 2 "omplower" } }
-! { dg-final { scan-tree-dump-times "private\\(k\\)" 1 "omplower" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/private-2.f95 b/gcc/testsuite/gfortran.dg/goacc/private-2.f95
deleted file mode 100644
index 4b038f2b5f20..000000000000
--- a/gcc/testsuite/gfortran.dg/goacc/private-2.f95
+++ /dev/null
@@ -1,39 +0,0 @@ 
-! { dg-do compile }
-
-! test for implicit private clauses in do loops
-
-program test
-  implicit none
-  integer :: i, j, k, a(10)
-
-  !$acc parallel
-  !$acc loop
-  do i = 1, 100
-  end do
-  !$acc end parallel
-
-  !$acc parallel
-  !$acc loop
-  do i = 1, 100
-     do j = 1, 100
-     end do
-  end do
-  !$acc end parallel
-
-  !$acc data copy(a)
-
-  if(mod(1,10) .eq. 0) write(*,'(i5)') i
-
-  do i = 1, 100
-    !$acc parallel
-    !$acc loop
-     do j = 1, 100
-        do k = 1, 100
-        end do
-     end do
-    !$acc end parallel
-  end do
-
-  !$acc end data
-
-end program test
diff --git a/gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f95 b/gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f95
new file mode 100644
index 000000000000..12a7854526a9
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f95
@@ -0,0 +1,248 @@ 
+! Predetermined 'private' clauses related to 'do' loops inside an OpenACC
+! 'kernels' construct.
+
+! { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" }
+
+! (The 'independent' clauses are used as end of directive markers in tree dump
+! scanning.)
+
+program test
+  implicit none
+  integer :: i0_1
+  integer :: i0_2, j0_2
+  integer :: i1_s
+  integer :: i1_c
+  integer :: i2_1_s, j2_1_s
+  integer :: i2_1_c, j2_1_c
+  integer :: i2_2_s, j2_2_s
+  integer :: i2_3_s, j2_3_s
+  integer :: i2_3_c, j2_3_c
+  integer :: i3_1_s, j3_1_s, k3_1_s
+  integer :: i3_1_c, j3_1_c, k3_1_c
+  integer :: i3_2_s, j3_2_s, k3_2_s
+  integer :: i3_2_c, j3_2_c, k3_2_c
+  integer :: i3_3_s, j3_3_s, k3_3_s
+  integer :: i3_3_c, j3_3_c, k3_3_c
+  integer :: i3_4_s, j3_4_s, k3_4_s
+  integer :: i3_4_c, j3_4_c, k3_4_c
+  integer :: i3_5_s, j3_5_s, k3_5_s
+
+  !$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
+  do i0_1 = 1, 100
+  end do
+  !$acc end kernels
+
+  !$acc kernels
+  ! { dg-final { scan-tree-dump-times "private\\(i0_2\\)" 1 "original" { xfail *-*-* } } } ! PR90067
+  ! { 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
+  do i0_2 = 1, 100
+     do j0_2 = 1, 100
+     end do
+  end do
+  !$acc end kernels
+
+  !$acc kernels
+  !$acc loop independent
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i1_s\\) independent" 1 "original" } }
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i1_s\\) independent" 1 "gimple" } }
+  do i1_s = 1, 100
+  end do
+  !$acc end kernels
+
+  !$acc kernels loop independent
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i1_c\\) independent" 1 "original" } }
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i1_c\\) independent" 1 "gimple" } }
+  do i1_c = 1, 100
+  end do
+
+  !$acc kernels
+  !$acc loop independent
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_s\\) private\\(j2_1_s\\) independent" 1 "original" } }
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_s\\) private\\(j2_1_s\\) independent" 1 "gimple" } }
+  do i2_1_s = 1, 100
+     do j2_1_s = 1, 100
+     end do
+  end do
+  !$acc end kernels
+
+  !$acc kernels loop independent
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_c\\) private\\(j2_1_c\\) independent" 1 "original" } }
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_c\\) private\\(j2_1_c\\) independent" 1 "gimple" } }
+  do i2_1_c = 1, 100
+     do j2_1_c = 1, 100
+     end do
+  end do
+
+  !$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
+  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" } }
+     ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_2_s\\) independent" 1 "gimple" } }
+     do j2_2_s = 1, 100
+     end do
+  end do
+  !$acc end kernels
+
+  !$acc kernels
+  !$acc loop independent
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_3_s\\) independent" 1 "original" } }
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_3_s\\) independent" 1 "gimple" } }
+  do i2_3_s = 1, 100
+     !$acc loop independent
+     ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_3_s\\) independent" 1 "original" } }
+     ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_3_s\\) independent" 1 "gimple" } }
+     do j2_3_s = 1, 100
+     end do
+  end do
+  !$acc end kernels
+
+  !$acc kernels loop independent
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_3_c\\) independent" 1 "original" } }
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_3_c\\) independent" 1 "gimple" } }
+  do i2_3_c = 1, 100
+     !$acc loop independent
+     ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_3_c\\) independent" 1 "original" } }
+     ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_3_c\\) independent" 1 "gimple" } }
+     do j2_3_c = 1, 100
+     end do
+  end do
+
+  !$acc kernels
+  !$acc loop independent
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_s\\) private\\(j3_1_s\\) private\\(k3_1_s\\) independent" 1 "original" } }
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_s\\) private\\(j3_1_s\\) private\\(k3_1_s\\) independent" 1 "gimple" } }
+  do i3_1_s = 1, 100
+     do j3_1_s = 1, 100
+        do k3_1_s = 1, 100
+        end do
+     end do
+  end do
+  !$acc end kernels
+
+  !$acc kernels loop independent
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_c\\) private\\(j3_1_c\\) private\\(k3_1_c\\) independent" 1 "original" } }
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_c\\) private\\(j3_1_c\\) private\\(k3_1_c\\) independent" 1 "gimple" } }
+  do i3_1_c = 1, 100
+     do j3_1_c = 1, 100
+        do k3_1_c = 1, 100
+        end do
+     end do
+  end do
+
+  !$acc kernels
+  !$acc loop independent
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_2_s\\) independent" 1 "original" } }
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_2_s\\) independent" 1 "gimple" } }
+  do i3_2_s = 1, 100
+     !$acc loop independent
+     ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_2_s\\) private\\(k3_2_s\\) independent" 1 "original" } }
+     ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_2_s\\) private\\(k3_2_s\\) independent" 1 "gimple" } }
+     do j3_2_s = 1, 100
+        do k3_2_s = 1, 100
+        end do
+     end do
+  end do
+  !$acc end kernels
+
+  !$acc kernels loop independent
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_2_c\\) independent" 1 "original" } }
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_2_c\\) independent" 1 "gimple" } }
+  do i3_2_c = 1, 100
+     !$acc loop independent
+     ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_2_c\\) private\\(k3_2_c\\) independent" 1 "original" } }
+     ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_2_c\\) private\\(k3_2_c\\) independent" 1 "gimple" } }
+     do j3_2_c = 1, 100
+        do k3_2_c = 1, 100
+        end do
+     end do
+  end do
+
+  !$acc kernels
+  !$acc loop independent
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_3_s\\) private\\(j3_3_s\\) independent" 1 "original" } }
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_3_s\\) private\\(j3_3_s\\) independent" 1 "gimple" } }
+  do i3_3_s = 1, 100
+     do j3_3_s = 1, 100
+        !$acc loop independent
+        ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_3_s\\) independent" 1 "original" } }
+        ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_3_s\\) independent" 1 "gimple" } }
+        do k3_3_s = 1, 100
+        end do
+     end do
+  end do
+  !$acc end kernels
+
+  !$acc kernels loop independent
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_3_c\\) private\\(j3_3_c\\) independent" 1 "original" } }
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_3_c\\) private\\(j3_3_c\\) independent" 1 "gimple" } }
+  do i3_3_c = 1, 100
+     do j3_3_c = 1, 100
+        !$acc loop independent
+        ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_3_c\\) independent" 1 "original" } }
+        ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_3_c\\) independent" 1 "gimple" } }
+        do k3_3_c = 1, 100
+        end do
+     end do
+  end do
+
+  !$acc kernels
+  !$acc loop independent
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_4_s\\) independent" 1 "original" } }
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_4_s\\) independent" 1 "gimple" } }
+  do i3_4_s = 1, 100
+     !$acc loop independent
+     ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_4_s\\) independent" 1 "original" } }
+     ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_4_s\\) independent" 1 "gimple" } }
+     do j3_4_s = 1, 100
+        !$acc loop independent
+        ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_4_s\\) independent" 1 "original" } }
+        ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_4_s\\) independent" 1 "gimple" } }
+        do k3_4_s = 1, 100
+        end do
+     end do
+  end do
+  !$acc end kernels
+
+  !$acc kernels loop independent
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_4_c\\) independent" 1 "original" } }
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_4_c\\) independent" 1 "gimple" } }
+  do i3_4_c = 1, 100
+     !$acc loop independent
+     ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_4_c\\) independent" 1 "original" } }
+     ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_4_c\\) independent" 1 "gimple" } }
+     do j3_4_c = 1, 100
+        !$acc loop independent
+        ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_4_c\\) independent" 1 "original" } }
+        ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_4_c\\) independent" 1 "gimple" } }
+        do k3_4_c = 1, 100
+        end do
+     end do
+  end do
+
+  !$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
+  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" } }
+     ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_5_s\\) independent" 1 "gimple" } }
+     do j3_5_s = 1, 100
+        !$acc loop independent
+        ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_5_s\\) independent" 1 "original" } }
+        ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_5_s\\) independent" 1 "gimple" } }
+        do k3_5_s = 1, 100
+        end do
+     end do
+  end do
+  !$acc end kernels
+end program test
diff --git a/gcc/testsuite/gfortran.dg/goacc/private-predetermined-parallel-1.f95 b/gcc/testsuite/gfortran.dg/goacc/private-predetermined-parallel-1.f95
new file mode 100644
index 000000000000..a2fa71a58b72
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/private-predetermined-parallel-1.f95
@@ -0,0 +1,253 @@ 
+! Predetermined 'private' clauses related to 'do' loops inside an OpenACC
+! 'parallel' construct.
+
+! { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" }
+
+! (The 'independent' clauses are used as end of directive markers in tree dump
+! scanning.)
+
+program test
+  implicit none
+  integer :: i0_1
+  integer :: i0_2, j0_2
+  integer :: i1_s
+  integer :: i1_c
+  integer :: i2_1_s, j2_1_s
+  integer :: i2_1_c, j2_1_c
+  integer :: i2_2_s, j2_2_s
+  integer :: i2_3_s, j2_3_s
+  integer :: i2_3_c, j2_3_c
+  integer :: i3_1_s, j3_1_s, k3_1_s
+  integer :: i3_1_c, j3_1_c, k3_1_c
+  integer :: i3_2_s, j3_2_s, k3_2_s
+  integer :: i3_2_c, j3_2_c, k3_2_c
+  integer :: i3_3_s, j3_3_s, k3_3_s
+  integer :: i3_3_c, j3_3_c, k3_3_c
+  integer :: i3_4_s, j3_4_s, k3_4_s
+  integer :: i3_4_c, j3_4_c, k3_4_c
+  integer :: i3_5_s, j3_5_s, k3_5_s
+
+  !$acc parallel
+  ! { dg-final { scan-tree-dump-times "private\\(i0_1\\)" 1 "original" { xfail *-*-* } } } ! PR90067
+  ! { dg-final { scan-tree-dump-times "private\\(i0_1\\)" 1 "gimple" } }
+  ! { dg-final { scan-tree-dump-times "firstprivate\\(i0_1\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+  ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_parallel firstprivate\\(i0_1\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+  do i0_1 = 1, 100
+  end do
+  !$acc end parallel
+
+  !$acc parallel
+  ! { dg-final { scan-tree-dump-times "private\\(i0_2\\)" 1 "original" { xfail *-*-* } } } ! PR90067
+  ! { dg-final { scan-tree-dump-times "private\\(j0_2\\)" 1 "original" { xfail *-*-* } } } ! PR90067
+  ! { dg-final { scan-tree-dump-times "private\\(i0_2\\)" 1 "gimple" } }
+  ! { dg-final { scan-tree-dump-times "private\\(j0_2\\)" 1 "gimple" } }
+  ! { dg-final { scan-tree-dump-times "firstprivate\\(i0_2\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+  ! { dg-final { scan-tree-dump-times "firstprivate\\(j0_2\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+  ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_parallel firstprivate\\(j0_2\\) firstprivate\\(i0_2\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+  do i0_2 = 1, 100
+     do j0_2 = 1, 100
+     end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel
+  !$acc loop independent
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i1_s\\) independent" 1 "original" } }
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i1_s\\) independent" 1 "gimple" } }
+  do i1_s = 1, 100
+  end do
+  !$acc end parallel
+
+  !$acc parallel loop independent
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i1_c\\) independent" 1 "original" } }
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i1_c\\) independent" 1 "gimple" } }
+  do i1_c = 1, 100
+  end do
+
+  !$acc parallel
+  !$acc loop independent
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_s\\) private\\(j2_1_s\\) independent" 1 "original" } }
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_s\\) private\\(j2_1_s\\) independent" 1 "gimple" } }
+  do i2_1_s = 1, 100
+     do j2_1_s = 1, 100
+     end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel loop independent
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_c\\) private\\(j2_1_c\\) independent" 1 "original" } }
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_c\\) private\\(j2_1_c\\) independent" 1 "gimple" } }
+  do i2_1_c = 1, 100
+     do j2_1_c = 1, 100
+     end do
+  end do
+
+  !$acc parallel
+  ! { 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" } }
+  ! { dg-final { scan-tree-dump-times "firstprivate\\(i2_2_s\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+  ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_parallel firstprivate\\(i2_2_s\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+  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" } }
+     ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_2_s\\) independent" 1 "gimple" } }
+     do j2_2_s = 1, 100
+     end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel
+  !$acc loop independent
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_3_s\\) independent" 1 "original" } }
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_3_s\\) independent" 1 "gimple" } }
+  do i2_3_s = 1, 100
+     !$acc loop independent
+     ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_3_s\\) independent" 1 "original" } }
+     ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_3_s\\) independent" 1 "gimple" } }
+     do j2_3_s = 1, 100
+     end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel loop independent
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_3_c\\) independent" 1 "original" } }
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_3_c\\) independent" 1 "gimple" } }
+  do i2_3_c = 1, 100
+     !$acc loop independent
+     ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_3_c\\) independent" 1 "original" } }
+     ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_3_c\\) independent" 1 "gimple" } }
+     do j2_3_c = 1, 100
+     end do
+  end do
+
+  !$acc parallel
+  !$acc loop independent
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_s\\) private\\(j3_1_s\\) private\\(k3_1_s\\) independent" 1 "original" } }
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_s\\) private\\(j3_1_s\\) private\\(k3_1_s\\) independent" 1 "gimple" } }
+  do i3_1_s = 1, 100
+     do j3_1_s = 1, 100
+        do k3_1_s = 1, 100
+        end do
+     end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel loop independent
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_c\\) private\\(j3_1_c\\) private\\(k3_1_c\\) independent" 1 "original" } }
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_c\\) private\\(j3_1_c\\) private\\(k3_1_c\\) independent" 1 "gimple" } }
+  do i3_1_c = 1, 100
+     do j3_1_c = 1, 100
+        do k3_1_c = 1, 100
+        end do
+     end do
+  end do
+
+  !$acc parallel
+  !$acc loop independent
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_2_s\\) independent" 1 "original" } }
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_2_s\\) independent" 1 "gimple" } }
+  do i3_2_s = 1, 100
+     !$acc loop independent
+     ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_2_s\\) private\\(k3_2_s\\) independent" 1 "original" } }
+     ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_2_s\\) private\\(k3_2_s\\) independent" 1 "gimple" } }
+     do j3_2_s = 1, 100
+        do k3_2_s = 1, 100
+        end do
+     end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel loop independent
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_2_c\\) independent" 1 "original" } }
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_2_c\\) independent" 1 "gimple" } }
+  do i3_2_c = 1, 100
+     !$acc loop independent
+     ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_2_c\\) private\\(k3_2_c\\) independent" 1 "original" } }
+     ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_2_c\\) private\\(k3_2_c\\) independent" 1 "gimple" } }
+     do j3_2_c = 1, 100
+        do k3_2_c = 1, 100
+        end do
+     end do
+  end do
+
+  !$acc parallel
+  !$acc loop independent
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_3_s\\) private\\(j3_3_s\\) independent" 1 "original" } }
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_3_s\\) private\\(j3_3_s\\) independent" 1 "gimple" } }
+  do i3_3_s = 1, 100
+     do j3_3_s = 1, 100
+        !$acc loop independent
+        ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_3_s\\) independent" 1 "original" } }
+        ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_3_s\\) independent" 1 "gimple" } }
+        do k3_3_s = 1, 100
+        end do
+     end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel loop independent
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_3_c\\) private\\(j3_3_c\\) independent" 1 "original" } }
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_3_c\\) private\\(j3_3_c\\) independent" 1 "gimple" } }
+  do i3_3_c = 1, 100
+     do j3_3_c = 1, 100
+        !$acc loop independent
+        ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_3_c\\) independent" 1 "original" } }
+        ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_3_c\\) independent" 1 "gimple" } }
+        do k3_3_c = 1, 100
+        end do
+     end do
+  end do
+
+  !$acc parallel
+  !$acc loop independent
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_4_s\\) independent" 1 "original" } }
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_4_s\\) independent" 1 "gimple" } }
+  do i3_4_s = 1, 100
+     !$acc loop independent
+     ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_4_s\\) independent" 1 "original" } }
+     ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_4_s\\) independent" 1 "gimple" } }
+     do j3_4_s = 1, 100
+        !$acc loop independent
+        ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_4_s\\) independent" 1 "original" } }
+        ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_4_s\\) independent" 1 "gimple" } }
+        do k3_4_s = 1, 100
+        end do
+     end do
+  end do
+  !$acc end parallel
+
+  !$acc parallel loop independent
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_4_c\\) independent" 1 "original" } }
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_4_c\\) independent" 1 "gimple" } }
+  do i3_4_c = 1, 100
+     !$acc loop independent
+     ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_4_c\\) independent" 1 "original" } }
+     ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_4_c\\) independent" 1 "gimple" } }
+     do j3_4_c = 1, 100
+        !$acc loop independent
+        ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_4_c\\) independent" 1 "original" } }
+        ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_4_c\\) independent" 1 "gimple" } }
+        do k3_4_c = 1, 100
+        end do
+     end do
+  end do
+
+  !$acc parallel
+  ! { 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" } }
+  ! { dg-final { scan-tree-dump-times "firstprivate\\(i3_5_s\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+  ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_parallel firstprivate\\(i3_5_s\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+  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" } }
+     ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_5_s\\) independent" 1 "gimple" } }
+     do j3_5_s = 1, 100
+        !$acc loop independent
+        ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_5_s\\) independent" 1 "original" } }
+        ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_5_s\\) independent" 1 "gimple" } }
+        do k3_5_s = 1, 100
+        end do
+     end do
+  end do
+  !$acc end parallel
+end program test
diff --git a/gcc/testsuite/gfortran.dg/goacc/private-predetermined-routine-1.f95 b/gcc/testsuite/gfortran.dg/goacc/private-predetermined-routine-1.f95
new file mode 100644
index 000000000000..6190669d63ae
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/private-predetermined-routine-1.f95
@@ -0,0 +1,142 @@ 
+! Predetermined 'private' clauses related to 'do' loops inside an OpenACC
+! accelerator routine.
+
+! { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" }
+
+! (The 'independent' clauses are used as end of directive markers in tree dump
+! scanning.)
+
+! The PR90114 XFAILs need to scan for the appropriate predetermined private
+! level.
+
+subroutine test
+  implicit none
+  integer :: i0_1
+  integer :: i0_2, j0_2
+  integer :: i1_s
+  integer :: i2_1_s, j2_1_s
+  integer :: i2_2_s, j2_2_s
+  integer :: i2_3_s, j2_3_s
+  integer :: i3_1_s, j3_1_s, k3_1_s
+  integer :: i3_2_s, j3_2_s, k3_2_s
+  integer :: i3_3_s, j3_3_s, k3_3_s
+  integer :: i3_4_s, j3_4_s, k3_4_s
+  integer :: i3_5_s, j3_5_s, k3_5_s
+  !$acc routine gang
+
+  ! { dg-final { scan-tree-dump-times "private\\(i0_1\\)" 1 "original" { xfail *-*-* } } } ! PR90114
+  ! { dg-final { scan-tree-dump-times "private\\(i0_1\\)" 1 "gimple" { xfail *-*-* } } } ! PR90114
+  do i0_1 = 1, 100
+  end do
+
+  ! { dg-final { scan-tree-dump-times "private\\(i0_2\\)" 1 "original" { xfail *-*-* } } } ! PR90114
+  ! { dg-final { scan-tree-dump-times "private\\(j0_2\\)" 1 "original" { xfail *-*-* } } } ! PR90114
+  ! { dg-final { scan-tree-dump-times "private\\(i0_2\\)" 1 "gimple" { xfail *-*-* } } } ! PR90114
+  ! { dg-final { scan-tree-dump-times "private\\(j0_2\\)" 1 "gimple" { xfail *-*-* } } } ! PR90114
+  do i0_2 = 1, 100
+     do j0_2 = 1, 100
+     end do
+  end do
+
+  !$acc loop independent
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i1_s\\) independent" 1 "original" } }
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i1_s\\) independent" 1 "gimple" } }
+  do i1_s = 1, 100
+  end do
+
+  !$acc loop independent
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_s\\) private\\(j2_1_s\\) independent" 1 "original" } }
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_s\\) private\\(j2_1_s\\) independent" 1 "gimple" } }
+  do i2_1_s = 1, 100
+     do j2_1_s = 1, 100
+     end do
+  end do
+
+  ! { dg-final { scan-tree-dump-times "private\\(i2_2_s\\)" 1 "original" { xfail *-*-* } } } ! PR90114
+  ! { dg-final { scan-tree-dump-times "private\\(i2_2_s\\)" 1 "gimple" { xfail *-*-* } } } ! PR90114
+  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" } }
+     ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_2_s\\) independent" 1 "gimple" } }
+     do j2_2_s = 1, 100
+     end do
+  end do
+
+  !$acc loop independent
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_3_s\\) independent" 1 "original" } }
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_3_s\\) independent" 1 "gimple" } }
+  do i2_3_s = 1, 100
+     !$acc loop independent
+     ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_3_s\\) independent" 1 "original" } }
+     ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_3_s\\) independent" 1 "gimple" } }
+     do j2_3_s = 1, 100
+     end do
+  end do
+
+  !$acc loop independent
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_s\\) private\\(j3_1_s\\) private\\(k3_1_s\\) independent" 1 "original" } }
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_s\\) private\\(j3_1_s\\) private\\(k3_1_s\\) independent" 1 "gimple" } }
+  do i3_1_s = 1, 100
+     do j3_1_s = 1, 100
+        do k3_1_s = 1, 100
+        end do
+     end do
+  end do
+
+  !$acc loop independent
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_2_s\\) independent" 1 "original" } }
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_2_s\\) independent" 1 "gimple" } }
+  do i3_2_s = 1, 100
+     !$acc loop independent
+     ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_2_s\\) private\\(k3_2_s\\) independent" 1 "original" } }
+     ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_2_s\\) private\\(k3_2_s\\) independent" 1 "gimple" } }
+     do j3_2_s = 1, 100
+        do k3_2_s = 1, 100
+        end do
+     end do
+  end do
+
+  !$acc loop independent
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_3_s\\) private\\(j3_3_s\\) independent" 1 "original" } }
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_3_s\\) private\\(j3_3_s\\) independent" 1 "gimple" } }
+  do i3_3_s = 1, 100
+     do j3_3_s = 1, 100
+        !$acc loop independent
+        ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_3_s\\) independent" 1 "original" } }
+        ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_3_s\\) independent" 1 "gimple" } }
+        do k3_3_s = 1, 100
+        end do
+     end do
+  end do
+
+  !$acc loop independent
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_4_s\\) independent" 1 "original" } }
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_4_s\\) independent" 1 "gimple" } }
+  do i3_4_s = 1, 100
+     !$acc loop independent
+     ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_4_s\\) independent" 1 "original" } }
+     ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_4_s\\) independent" 1 "gimple" } }
+     do j3_4_s = 1, 100
+        !$acc loop independent
+        ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_4_s\\) independent" 1 "original" } }
+        ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_4_s\\) independent" 1 "gimple" } }
+        do k3_4_s = 1, 100
+        end do
+     end do
+  end do
+
+  ! { dg-final { scan-tree-dump-times "private\\(i3_5_s\\)" 1 "original" { xfail *-*-* } } } ! PR90114
+  ! { dg-final { scan-tree-dump-times "private\\(i3_5_s\\)" 1 "gimple" { xfail *-*-* } } } ! PR90114
+  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" } }
+     ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_5_s\\) independent" 1 "gimple" } }
+     do j3_5_s = 1, 100
+        !$acc loop independent
+        ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_5_s\\) independent" 1 "original" } }
+        ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(k3_5_s\\) independent" 1 "gimple" } }
+        do k3_5_s = 1, 100
+        end do
+     end do
+  end do
+end subroutine test
-- 
2.17.1