[PR middle-end/71373] Handle more OMP_CLAUSE_* in nested function decomposition

Backport trunk r237291:

	gcc/
	* gimplify.c (gimplify_adjust_omp_clauses): Discard
	OMP_CLAUSE_TILE.
	* omp-low.c (scan_sharing_clauses): Don't expect OMP_CLAUSE_TILE.
	gcc/testsuite/
	* c-c++-common/goacc/combined-directives.c: XFAIL tree scanning
	for OpenACC tile clauses.
	* gfortran.dg/goacc/combined-directives.f90: Likewise.

	gcc/
	PR middle-end/71373
	* tree-nested.c (convert_nonlocal_omp_clauses)
	(convert_local_omp_clauses): Handle OMP_CLAUSE_ASYNC,
	OMP_CLAUSE_WAIT, OMP_CLAUSE_INDEPENDENT, OMP_CLAUSE_AUTO,
	OMP_CLAUSE__CACHE_, OMP_CLAUSE_TILE.
	gcc/testsuite/
	PR middle-end/71373
	* gcc.dg/goacc/nested-function-1.c: New file.
	* gcc.dg/goacc/nested-function-2.c: Likewise.
	* gcc.dg/goacc/pr71373.c: Likewise.
	* gfortran.dg/goacc/cray-2.f95: Likewise.
	* gfortran.dg/goacc/loop-1-2.f95: Likewise.
	* gfortran.dg/goacc/loop-3-2.f95: Likewise.
	* gfortran.dg/goacc/cray.f95: Update.
	* gfortran.dg/goacc/loop-1.f95: Likewise.
	* gfortran.dg/goacc/loop-3.f95: Likewise.
	* gfortran.dg/goacc/subroutines.f90: Update, and rename to...
	* gfortran.dg/goacc/nested-function-1.f90: ... this new file.
	libgomp/testsuite/
	PR middle-end/71373
	* libgomp.oacc-c/nested-function-1.c: New file.
	* libgomp.oacc-c/nested-function-2.c: Likewise.
	* libgomp.oacc-fortran/nested-function-1.f90: Likewise.
	* libgomp.oacc-fortran/nested-function-2.f90: Likewise.
	* libgomp.oacc-fortran/nested-function-3.f90: Likewise.

git-svn-id: https://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@237300 138bc75d-0d04-0410-961f-82ee72b054a4

18 files changed, 679 insertions, 100 deletions

diff --git a/gcc/ChangeLog.gomp b/gcc/ChangeLog.gomp
index 4477abfa462..b68538d6034 100644
--- a/gcc/ChangeLog.gomp
+++ b/gcc/ChangeLog.gomp
@@ -1,5 +1,17 @@
 2016-06-10  Thomas Schwinge  <thomas@codesourcery.com>
 
+	PR middle-end/71373
+	Backport from trunk r237291:
+	* tree-nested.c (convert_nonlocal_omp_clauses)
+	(convert_local_omp_clauses): Handle OMP_CLAUSE_ASYNC,
+	OMP_CLAUSE_WAIT, OMP_CLAUSE_INDEPENDENT, OMP_CLAUSE_AUTO,
+	OMP_CLAUSE__CACHE_, OMP_CLAUSE_TILE.
+
+	Backport from trunk r237291:
+	* gimplify.c (gimplify_adjust_omp_clauses): Discard
+	OMP_CLAUSE_TILE.
+	* omp-low.c (scan_sharing_clauses): Don't expect OMP_CLAUSE_TILE.
+
 	Backport from trunk r237290:
 	* omp-low.c (scan_sharing_clauses): Don't expect
 	OMP_CLAUSE__CACHE_.
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 37971c7085e..717b25f5fb1 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -8265,10 +8265,16 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
 	case OMP_CLAUSE_VECTOR:
 	case OMP_CLAUSE_AUTO:
 	case OMP_CLAUSE_SEQ:
-	case OMP_CLAUSE_TILE:
 	case OMP_CLAUSE_DEVICE_TYPE:
 	  break;
 
+	case OMP_CLAUSE_TILE:
+	  /* We're not yet making use of the information provided by OpenACC
+	     tile clauses.  Discard these here, to simplify later middle end
+	     processing.  */
+	  remove = true;
+	  break;
+
 	case OMP_CLAUSE_BIND:
 	case OMP_CLAUSE_NOHOST:
 	default:
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 40ac8c86ec2..fedb1958db1 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -2218,7 +2218,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx,
 	case OMP_CLAUSE_GANG:
 	case OMP_CLAUSE_WORKER:
 	case OMP_CLAUSE_VECTOR:
-	case OMP_CLAUSE_TILE:
 	case OMP_CLAUSE_INDEPENDENT:
 	case OMP_CLAUSE_AUTO:
 	case OMP_CLAUSE_SEQ:
@@ -2235,6 +2234,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx,
 	case OMP_CLAUSE_BIND:
 	case OMP_CLAUSE_DEVICE_RESIDENT:
 	case OMP_CLAUSE_NOHOST:
+	case OMP_CLAUSE_TILE:
 	case OMP_CLAUSE__CACHE_:
 	default:
 	  gcc_unreachable ();
@@ -2392,7 +2392,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx,
 	case OMP_CLAUSE_GANG:
 	case OMP_CLAUSE_WORKER:
 	case OMP_CLAUSE_VECTOR:
-	case OMP_CLAUSE_TILE:
 	case OMP_CLAUSE_INDEPENDENT:
 	case OMP_CLAUSE_AUTO:
 	case OMP_CLAUSE_SEQ:
@@ -2403,6 +2402,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx,
 	case OMP_CLAUSE_BIND:
 	case OMP_CLAUSE_DEVICE_RESIDENT:
 	case OMP_CLAUSE_NOHOST:
+	case OMP_CLAUSE_TILE:
 	case OMP_CLAUSE__CACHE_:
 	default:
 	  gcc_unreachable ();
diff --git a/gcc/testsuite/ChangeLog.gomp b/gcc/testsuite/ChangeLog.gomp
index eef9425ea2f..031774da7ad 100644
--- a/gcc/testsuite/ChangeLog.gomp
+++ b/gcc/testsuite/ChangeLog.gomp
@@ -1,5 +1,27 @@
 2016-06-10  Thomas Schwinge  <thomas@codesourcery.com>
 
+	PR middle-end/71373
+	Backport from trunk r237291:
+	2016-06-10  Thomas Schwinge  <thomas@codesourcery.com>
+		    Cesar Philippidis  <cesar@codesourcery.com>
+
+	* gcc.dg/goacc/nested-function-1.c: New file.
+	* gcc.dg/goacc/nested-function-2.c: Likewise.
+	* gcc.dg/goacc/pr71373.c: Likewise.
+	* gfortran.dg/goacc/cray-2.f95: Likewise.
+	* gfortran.dg/goacc/loop-1-2.f95: Likewise.
+	* gfortran.dg/goacc/loop-3-2.f95: Likewise.
+	* gfortran.dg/goacc/cray.f95: Update.
+	* gfortran.dg/goacc/loop-1.f95: Likewise.
+	* gfortran.dg/goacc/loop-3.f95: Likewise.
+	* gfortran.dg/goacc/subroutines.f90: Update, and rename to...
+	* gfortran.dg/goacc/nested-function-1.f90: ... this new file.
+
+	Backport from trunk r237291:
+	* c-c++-common/goacc/combined-directives.c: XFAIL tree scanning
+	for OpenACC tile clauses.
+	* gfortran.dg/goacc/combined-directives.f90: Likewise.
+
 	PR c/71381
 	Backport from trunk r237290:
 	* c-c++-common/goacc/cache-1.c: Update.  Move invalid usage tests
diff --git a/gcc/testsuite/c-c++-common/goacc/combined-directives.c b/gcc/testsuite/c-c++-common/goacc/combined-directives.c
index 2ef5b537717..2342c9fb63f 100644
--- a/gcc/testsuite/c-c++-common/goacc/combined-directives.c
+++ b/gcc/testsuite/c-c++-common/goacc/combined-directives.c
@@ -114,6 +114,7 @@ test ()
 // { dg-final { scan-tree-dump-times "acc loop vector" 2 "gimple" } }
 // { dg-final { scan-tree-dump-times "acc loop seq" 2 "gimple" } }
 // { dg-final { scan-tree-dump-times "acc loop auto" 2 "gimple" } }
-// { dg-final { scan-tree-dump-times "acc loop tile.2, 3" 2 "gimple" } }
+// XFAILed: OpenACC tile clauses are discarded during gimplification.
+// { dg-final { scan-tree-dump-times "acc loop tile.2, 3" 2 "gimple" { xfail *-*-* } } }
 // { dg-final { scan-tree-dump-times "acc loop independent private.i" 2 "gimple" } }
 // { dg-final { scan-tree-dump-times "private.z" 2 "gimple" } }
diff --git a/gcc/testsuite/gcc.dg/goacc/nested-function-1.c b/gcc/testsuite/gcc.dg/goacc/nested-function-1.c
new file mode 100644
index 00000000000..e17c0e2227f
--- /dev/null
+++ b/gcc/testsuite/gcc.dg/goacc/nested-function-1.c
@@ -0,0 +1,100 @@
+/* Exercise nested function decomposition, gcc/tree-nested.c.  */
+/* See gcc/testsuite/gfortran.dg/goacc/nested-function-1.f90 for the Fortran
+   version.  */
+
+int main ()
+{
+#define N 100
+  int nonlocal_arg;
+  int nonlocal_a[N];
+  int nonlocal_i;
+  int nonlocal_j;
+
+  for (int i = 0; i < N; ++i)
+    nonlocal_a[i] = 5;
+  nonlocal_arg = 5;
+
+  void local ()
+  {
+    int local_i;
+    int local_arg;
+    int local_a[N];
+    int local_j;
+
+    for (int i = 0; i < N; ++i)
+      local_a[i] = 5;
+    local_arg = 5;
+
+#pragma acc kernels loop \
+  gang(num:local_arg) worker(local_arg) vector(local_arg) \
+  wait async(local_arg)
+    for (local_i = 0; local_i < N; ++local_i)
+      {
+#pragma acc cache (local_a[local_i:5])
+	local_a[local_i] = 100;
+#pragma acc loop seq tile(*)
+	for (local_j = 0; local_j < N; ++local_j)
+	  ;
+#pragma acc loop auto independent tile(1)
+	for (local_j = 0; local_j < N; ++local_j)
+	  ;
+      }
+
+#pragma acc kernels loop \
+  gang(static:local_arg) worker(local_arg) vector(local_arg) \
+  wait(local_arg, local_arg + 1, local_arg + 2) async
+    for (local_i = 0; local_i < N; ++local_i)
+      {
+#pragma acc cache (local_a[local_i:4])
+	local_a[local_i] = 100;
+#pragma acc loop seq tile(1)
+	for (local_j = 0; local_j < N; ++local_j)
+	  ;
+#pragma acc loop auto independent tile(*)
+	for (local_j = 0; local_j < N; ++local_j)
+	  ;
+      }
+  }
+
+  void nonlocal ()
+  {
+    for (int i = 0; i < N; ++i)
+      nonlocal_a[i] = 5;
+    nonlocal_arg = 5;
+
+#pragma acc kernels loop \
+  gang(num:nonlocal_arg) worker(nonlocal_arg) vector(nonlocal_arg) \
+  wait async(nonlocal_arg)
+    for (nonlocal_i = 0; nonlocal_i < N; ++nonlocal_i)
+      {
+#pragma acc cache (nonlocal_a[nonlocal_i:3])
+	nonlocal_a[nonlocal_i] = 100;
+#pragma acc loop seq tile(2)
+	for (nonlocal_j = 0; nonlocal_j < N; ++nonlocal_j)
+	  ;
+#pragma acc loop auto independent tile(3)
+	for (nonlocal_j = 0; nonlocal_j < N; ++nonlocal_j)
+	  ;
+      }
+
+#pragma acc kernels loop \
+  gang(static:nonlocal_arg) worker(nonlocal_arg) vector(nonlocal_arg) \
+  wait(nonlocal_arg, nonlocal_arg + 1, nonlocal_arg + 2) async
+    for (nonlocal_i = 0; nonlocal_i < N; ++nonlocal_i)
+      {
+#pragma acc cache (nonlocal_a[nonlocal_i:2])
+	nonlocal_a[nonlocal_i] = 100;
+#pragma acc loop seq tile(*)
+	for (nonlocal_j = 0; nonlocal_j < N; ++nonlocal_j)
+	  ;
+#pragma acc loop auto independent tile(*)
+	for (nonlocal_j = 0; nonlocal_j < N; ++nonlocal_j)
+	  ;
+      }
+  }
+
+  local ();
+  nonlocal ();
+
+  return 0;
+}
diff --git a/gcc/testsuite/gcc.dg/goacc/nested-function-2.c b/gcc/testsuite/gcc.dg/goacc/nested-function-2.c
new file mode 100644
index 00000000000..70c9ec8ebfa
--- /dev/null
+++ b/gcc/testsuite/gcc.dg/goacc/nested-function-2.c
@@ -0,0 +1,45 @@
+/* Exercise nested function decomposition, gcc/tree-nested.c.  */
+
+int
+main (void)
+{
+  int j = 0, k = 6, l = 7, m = 8;
+  void simple (void)
+  {
+    int i;
+#pragma acc parallel
+    {
+#pragma acc loop
+      for (i = 0; i < m; i+= k)
+	j = (m + i - j) * l;
+    }
+  }
+  void collapse (void)
+  {
+    int x, y, z;
+#pragma acc parallel
+    {
+#pragma acc loop collapse (3)
+      for (x = 0; x < k; x++)
+	for (y = -5; y < l; y++)
+	  for (z = 0; z < m; z++)
+	    j += x + y + z;
+    }
+  }
+  void reduction (void)
+  {
+    int x, y, z;
+#pragma acc parallel reduction (+:j)
+    {
+#pragma acc loop reduction (+:j) collapse (3)
+      for (x = 0; x < k; x++)
+	for (y = -5; y < l; y++)
+	  for (z = 0; z < m; z++)
+	    j += x + y + z;
+    }
+  }
+  simple();
+  collapse();
+  reduction();
+  return 0;
+}
diff --git a/gcc/testsuite/gcc.dg/goacc/pr71373.c b/gcc/testsuite/gcc.dg/goacc/pr71373.c
new file mode 100644
index 00000000000..9381752cc9d
--- /dev/null
+++ b/gcc/testsuite/gcc.dg/goacc/pr71373.c
@@ -0,0 +1,41 @@
+/* Unintentional nested function usage.  */
+/* Due to missing right braces '}', the following functions are parsed as
+   nested functions.  This ran into an ICE.  */
+
+void foo (void)
+{
+  #pragma acc parallel
+  {
+    #pragma acc loop independent
+    for (int i = 0; i < 16; i++)
+      ;
+  // Note right brace '}' commented out here.
+  //}
+}
+void bar (void)
+{
+}
+
+// Adding right brace '}' here, to make this compile.
+}
+
+
+// ..., and the other way round:
+
+void BAR (void)
+{
+// Note right brace '}' commented out here.
+//}
+
+void FOO (void)
+{
+  #pragma acc parallel
+  {
+    #pragma acc loop independent
+    for (int i = 0; i < 16; i++)
+      ;
+  }
+}
+
+// Adding right brace '}' here, to make this compile.
+}
diff --git a/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90 b/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90
index 00b8822338c..ffa93715ca3 100644
--- a/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90
@@ -162,6 +162,7 @@ end subroutine test
 ! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. vector" 2 "gimple" { xfail *-*-* } } }
 ! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. seq" 2 "gimple" { xfail *-*-* } } }
 ! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. auto" 2 "gimple" { xfail *-*-* } } }
+! XFAILed: OpenACC tile clauses are discarded during gimplification.
 ! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. tile.2, 3" 2 "gimple" { xfail *-*-* } } }
 ! { dg-final { scan-tree-dump-times "acc loop private.i. independent" 2 "gimple" { xfail *-*-* } } }
 ! { dg-final { scan-tree-dump-times "private.z" 2 "gimple" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/cray-2.f95 b/gcc/testsuite/gfortran.dg/goacc/cray-2.f95
new file mode 100644
index 00000000000..51b79b53636
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/cray-2.f95
@@ -0,0 +1,56 @@
+! { dg-additional-options "-fcray-pointer" }
+! See also cray.f95.
+
+program test
+  call oacc1
+contains
+  subroutine oacc1
+    implicit none
+    integer :: i
+    real :: pointee
+    pointer (ptr, pointee)
+    !$acc declare device_resident (pointee)
+    !$acc declare device_resident (ptr)
+    !$acc data copy (pointee) ! { dg-error "Cray pointee" }
+    !$acc end data
+    !$acc data deviceptr (pointee) ! { dg-error "Cray pointee" }
+    !$acc end data
+    !$acc parallel private (pointee) ! { dg-error "Cray pointee" }
+    !$acc end parallel
+    !$acc host_data use_device (pointee) ! { dg-error "Cray pointee" }
+    !$acc end host_data
+    !$acc parallel loop reduction(+:pointee) ! { dg-error "Cray pointee" }
+    do i = 1,5
+    enddo
+    !$acc end parallel loop
+    !$acc parallel loop
+    do i = 1,5
+      !$acc cache (pointee) ! { dg-error "Cray pointee" }
+    enddo
+    !$acc end parallel loop
+    !$acc update device (pointee) ! { dg-error "Cray pointee" }
+    !$acc update host (pointee) ! { dg-error "Cray pointee" }
+    !$acc update self (pointee) ! { dg-error "Cray pointee" }
+    !$acc data copy (ptr)
+    !$acc end data
+    !$acc data deviceptr (ptr) ! { dg-error "Cray pointer" }
+    !$acc end data
+    !$acc parallel private (ptr)
+    !$acc end parallel
+    !$acc host_data use_device (ptr) ! { dg-error "Cray pointer" }
+    !$acc end host_data
+    !$acc parallel loop reduction(+:ptr) ! { dg-error "Cray pointer" }
+    do i = 1,5
+    enddo
+    !$acc end parallel loop
+    !$acc parallel loop
+    do i = 1,5
+      !TODO: This must fail, as in openacc-1_0-branch.
+      !$acc cache (ptr) ! { dg-error "" "TODO" { xfail *-*-* } }
+    enddo
+    !$acc end parallel loop
+    !$acc update device (ptr)
+    !$acc update host (ptr)
+    !$acc update self (ptr)
+  end subroutine oacc1
+end program test
diff --git a/gcc/testsuite/gfortran.dg/goacc/cray.f95 b/gcc/testsuite/gfortran.dg/goacc/cray.f95
index 705c18c992d..d6d531705a6 100644
--- a/gcc/testsuite/gfortran.dg/goacc/cray.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/cray.f95
@@ -1,5 +1,5 @@
-! { dg-do compile } 
 ! { dg-additional-options "-fcray-pointer" }
+! See also cray-2.f95.
 
 module test
 contains
@@ -8,8 +8,8 @@ contains
     integer :: i
     real :: pointee
     pointer (ptr, pointee)
-    !$acc declare device_resident (pointee) 
-    !$acc declare device_resident (ptr) 
+    !$acc declare device_resident (pointee)
+    !$acc declare device_resident (ptr)
     !$acc data copy (pointee) ! { dg-error "Cray pointee" }
     !$acc end data
     !$acc data deviceptr (pointee) ! { dg-error "Cray pointee" }
diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-1-2.f95 b/gcc/testsuite/gfortran.dg/goacc/loop-1-2.f95
new file mode 100644
index 00000000000..79665b948c3
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/loop-1-2.f95
@@ -0,0 +1,176 @@
+! See also loop-1.f95.
+
+program test
+  call test1
+contains
+
+subroutine test1
+  integer :: i, j, k, b(10)
+  integer, dimension (30) :: a
+  double precision :: d
+  real :: r
+  i = 0
+  !$acc loop
+  do 100 ! { dg-error "cannot be a DO WHILE or DO without loop control" }
+    if (i .gt. 0) exit ! { dg-error "EXIT statement" }
+  100 i = i + 1
+  i = 0
+  !$acc loop
+  do ! { dg-error "cannot be a DO WHILE or DO without loop control" }
+      if (i .gt. 0) exit ! { dg-error "EXIT statement" }
+       i = i + 1
+  end do
+  i = 0
+  !$acc loop
+  do 200 while (i .lt. 4) ! { dg-error "cannot be a DO WHILE or DO without loop control" }
+  200 i = i + 1
+  !$acc loop
+  do while (i .lt. 8) ! { dg-error "cannot be a DO WHILE or DO without loop control" }
+       i = i + 1
+  end do
+  !$acc loop
+  do 300 d = 1, 30, 6
+      i = d
+  300 a(i) = 1
+  ! { dg-warning "Deleted feature: Loop variable at .1. must be integer" "" { target *-*-* } 32 }
+  ! { dg-error "ACC LOOP iteration variable must be of type integer" "" { target *-*-* } 32 }
+  !$acc loop
+  do d = 1, 30, 5
+       i = d
+      a(i) = 2
+  end do
+  ! { dg-warning "Deleted feature: Loop variable at .1. must be integer" "" { target *-*-* } 38 }
+  ! { dg-error "ACC LOOP iteration variable must be of type integer" "" { target *-*-* } 38 }
+  !$acc loop
+  do i = 1, 30
+      if (i .eq. 16) exit ! { dg-error "EXIT statement" }
+  end do
+  !$acc loop
+  outer: do i = 1, 30
+      do j = 5, 10
+          if (i .eq. 6 .and. j .eq. 7) exit outer ! { dg-error "EXIT statement" }
+      end do
+  end do outer
+  last: do i = 1, 30
+   end do last
+
+  ! different types of loop are allowed
+  !$acc loop
+  do i = 1,10
+  end do
+  !$acc loop
+  do 400, i = 1,10
+400   a(i) = i
+
+  ! after loop directive must be loop
+  !$acc loop
+  a(1) = 1 ! { dg-error "Expected DO loop" }
+  do i = 1,10
+  enddo
+
+  ! combined directives may be used with/without end
+  !$acc parallel loop
+  do i = 1,10
+  enddo
+  !$acc parallel loop
+  do i = 1,10
+  enddo
+  !$acc end parallel loop
+  !$acc kernels loop
+  do i = 1,10
+  enddo
+  !$acc kernels loop
+  do i = 1,10
+  enddo
+  !$acc end kernels loop
+
+  !$acc kernels loop reduction(max:i)
+  do i = 1,10
+  enddo
+  !$acc kernels
+  !$acc loop reduction(max:i)
+  do i = 1,10
+  enddo
+  !$acc end kernels
+
+  !$acc parallel loop collapse(0) ! { dg-error "constant positive integer" }
+  do i = 1,10
+  enddo
+
+  !$acc parallel loop collapse(-1) ! { dg-error "constant positive integer" }
+  do i = 1,10
+  enddo
+
+  !$acc parallel loop collapse(i) ! { dg-error "Constant expression required" }
+  do i = 1,10
+  enddo
+
+  !$acc parallel loop collapse(4) ! { dg-error "not enough DO loops for collapsed" }
+    do i = 1, 3
+        do j = 4, 6
+          do k = 5, 7
+              a(i+j-k) = i + j + k
+          end do
+        end do
+    end do
+    !$acc parallel loop collapse(2)
+    do i = 1, 5, 2
+        do j = i + 1, 7, i  ! { dg-error "collapsed loops don.t form rectangular iteration space" }
+        end do
+    end do
+    !$acc parallel loop collapse(2)
+    do i = 1, 3
+        do j = 4, 6
+        end do
+    end do
+    !$acc parallel loop collapse(2)
+    do i = 1, 3
+        do j = 4, 6
+        end do
+        k = 4
+    end do
+    !$acc parallel loop collapse(3-1)
+    do i = 1, 3
+        do j = 4, 6
+        end do
+        k = 4
+    end do
+    !$acc parallel loop collapse(1+1)
+    do i = 1, 3
+        do j = 4, 6
+        end do
+        k = 4
+    end do
+    !$acc parallel loop collapse(2)
+    do i = 1, 3
+        do      ! { dg-error "cannot be a DO WHILE or DO without loop control" }
+        end do
+    end do
+    !$acc parallel loop collapse(2)
+    do i = 1, 3
+        do r = 4, 6
+        end do
+        ! { dg-warning "Deleted feature: Loop variable at .1. must be integer" "" { target *-*-* } 151 }
+        ! { dg-error "ACC LOOP iteration variable must be of type integer" "" { target *-*-* } 151 }
+    end do
+
+    ! Both seq and independent are not allowed
+  !$acc loop independent seq ! { dg-error "SEQ conflicts with INDEPENDENT" }
+  do i = 1,10
+  enddo
+
+
+  !$acc cache (a(1:10)) ! { dg-error "ACC CACHE directive must be inside of loop" }
+
+  do i = 1,10
+    !$acc cache(a(i:i+1))
+  enddo
+
+  do i = 1,10
+    !$acc cache(a(i:i+1))
+    a(i) = i
+    !$acc cache(a(i+2:i+2+1))
+  enddo
+
+end subroutine test1
+end program test
diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-1.f95 b/gcc/testsuite/gfortran.dg/goacc/loop-1.f95
index a605f038926..5f81b7a1d19 100644
--- a/gcc/testsuite/gfortran.dg/goacc/loop-1.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/loop-1.f95
@@ -1,8 +1,10 @@
+! See also loop-1-2.f95.
+
 module test
   implicit none
 contains
 
-subroutine test1  
+subroutine test1
   integer :: i, j, k, b(10)
   integer, dimension (30) :: a
   double precision :: d
@@ -30,15 +32,15 @@ subroutine test1
   do 300 d = 1, 30, 6
       i = d
   300 a(i) = 1
-  ! { dg-warning "Deleted feature: Loop variable at .1. must be integer" "" { target *-*-* } 30 }
-  ! { dg-error "ACC LOOP iteration variable must be of type integer" "" { target *-*-* } 30 }
+  ! { dg-warning "Deleted feature: Loop variable at .1. must be integer" "" { target *-*-* } 32 }
+  ! { dg-error "ACC LOOP iteration variable must be of type integer" "" { target *-*-* } 32 }
   !$acc loop
   do d = 1, 30, 5
        i = d
       a(i) = 2
   end do
-  ! { dg-warning "Deleted feature: Loop variable at .1. must be integer" "" { target *-*-* } 36 }
-  ! { dg-error "ACC LOOP iteration variable must be of type integer" "" { target *-*-* } 36 }
+  ! { dg-warning "Deleted feature: Loop variable at .1. must be integer" "" { target *-*-* } 38 }
+  ! { dg-error "ACC LOOP iteration variable must be of type integer" "" { target *-*-* } 38 }
   !$acc loop
   do i = 1, 30
       if (i .eq. 16) exit ! { dg-error "EXIT statement" }
@@ -53,7 +55,7 @@ subroutine test1
    end do last
 
   ! different types of loop are allowed
-  !$acc loop 
+  !$acc loop
   do i = 1,10
   end do
   !$acc loop
@@ -65,8 +67,8 @@ subroutine test1
   a(1) = 1 ! { dg-error "Expected DO loop" }
   do i = 1,10
   enddo
-  
-  ! combined directives may be used with/without end 
+
+  ! combined directives may be used with/without end
   !$acc parallel loop
   do i = 1,10
   enddo
@@ -82,11 +84,11 @@ subroutine test1
   enddo
   !$acc end kernels loop
 
-  !$acc kernels loop reduction(max:i) 
+  !$acc kernels loop reduction(max:i)
   do i = 1,10
   enddo
-  !$acc kernels 
-  !$acc loop reduction(max:i) 
+  !$acc kernels
+  !$acc loop reduction(max:i)
   do i = 1,10
   enddo
   !$acc end kernels
@@ -118,7 +120,7 @@ subroutine test1
     end do
     !$acc parallel loop collapse(2)
     do i = 1, 3
-        do j = 4, 6  
+        do j = 4, 6
         end do
     end do
     !$acc parallel loop collapse(2)
@@ -148,8 +150,8 @@ subroutine test1
     do i = 1, 3
         do r = 4, 6
         end do
-        ! { dg-warning "Deleted feature: Loop variable at .1. must be integer" "" { target *-*-* } 149 }
-        ! { dg-error "ACC LOOP iteration variable must be of type integer" "" { target *-*-* } 149 }
+        ! { dg-warning "Deleted feature: Loop variable at .1. must be integer" "" { target *-*-* } 151 }
+        ! { dg-error "ACC LOOP iteration variable must be of type integer" "" { target *-*-* } 151 }
     end do
 
     ! Both seq and independent are not allowed
diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-3-2.f95 b/gcc/testsuite/gfortran.dg/goacc/loop-3-2.f95
new file mode 100644
index 00000000000..9be74a85919
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/loop-3-2.f95
@@ -0,0 +1,58 @@
+! { dg-additional-options "-std=f2008" }
+! See also loop-3.f95.
+
+program test
+  call test1
+contains
+subroutine test1
+  implicit none
+  integer :: i, j
+
+  ! !$acc end loop not required by spec
+  !$acc loop
+  do i = 1,5
+  enddo
+  !$acc end loop ! { dg-warning "Redundant" }
+
+  !$acc loop
+  do i = 1,5
+  enddo
+  j = 1
+  !$acc end loop ! { dg-error "Unexpected" }
+
+  !$acc parallel
+  !$acc loop
+  do i = 1,5
+  enddo
+  !$acc end parallel
+  !$acc end loop ! { dg-error "Unexpected" }
+
+  ! OpenACC supports Fortran 2008 do concurrent statement
+  !$acc loop
+  do concurrent (i = 1:5)
+  end do
+
+  !$acc loop
+  outer_loop: do i = 1, 5
+    inner_loop: do j = 1,5
+      if (i .eq. j) cycle outer_loop
+      if (i .ne. j) exit outer_loop ! { dg-error "EXIT statement" }
+    end do inner_loop
+  end do outer_loop
+
+  outer_loop1: do i = 1, 5
+    !$acc loop
+    inner_loop1: do j = 1,5
+      if (i .eq. j) cycle outer_loop1 ! { dg-error "CYCLE statement" }
+    end do inner_loop1
+  end do outer_loop1
+
+  !$acc loop collapse(2)
+  outer_loop2: do i = 1, 5
+    inner_loop2: do j = 1,5
+      if (i .eq. j) cycle outer_loop2 ! { dg-error "CYCLE statement" }
+      if (i .ne. j) exit outer_loop2 ! { dg-error "EXIT statement" }
+    end do inner_loop2
+  end do outer_loop2
+end subroutine test1
+end program test
diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-3.f95 b/gcc/testsuite/gfortran.dg/goacc/loop-3.f95
index 2a866c79234..30930f404f3 100644
--- a/gcc/testsuite/gfortran.dg/goacc/loop-3.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/loop-3.f95
@@ -1,10 +1,10 @@
-! { dg-do compile }
 ! { dg-additional-options "-std=f2008" }
+! See also loop-3-2.f95.
 
 subroutine test1
   implicit none
   integer :: i, j
-  
+
   ! !$acc end loop not required by spec
   !$acc loop
   do i = 1,5
@@ -23,7 +23,7 @@ subroutine test1
   enddo
   !$acc end parallel
   !$acc end loop ! { dg-error "Unexpected" }
-  
+
   ! OpenACC supports Fortran 2008 do concurrent statement
   !$acc loop
   do concurrent (i = 1:5)
@@ -35,7 +35,7 @@ subroutine test1
       if (i .eq. j) cycle outer_loop
       if (i .ne. j) exit outer_loop ! { dg-error "EXIT statement" }
     end do inner_loop
-  end do outer_loop 
+  end do outer_loop
 
   outer_loop1: do i = 1, 5
     !$acc loop
@@ -50,6 +50,5 @@ subroutine test1
       if (i .eq. j) cycle outer_loop2 ! { dg-error "CYCLE statement" }
       if (i .ne. j) exit outer_loop2 ! { dg-error "EXIT statement" }
     end do inner_loop2
-  end do outer_loop2 
+  end do outer_loop2
 end subroutine test1
-
diff --git a/gcc/testsuite/gfortran.dg/goacc/nested-function-1.f90 b/gcc/testsuite/gfortran.dg/goacc/nested-function-1.f90
new file mode 100644
index 00000000000..2fcaa400ee3
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/nested-function-1.f90
@@ -0,0 +1,93 @@
+! Exercise nested function decomposition, gcc/tree-nested.c.
+! See gcc/testsuite/gcc.dg/goacc/nested-function-1.c for the C version.
+
+program main
+  integer, parameter :: N = 100
+  integer :: nonlocal_arg
+  integer :: nonlocal_a(N)
+  integer :: nonlocal_i
+  integer :: nonlocal_j
+
+  nonlocal_a (:) = 5
+  nonlocal_arg = 5
+
+  call local ()
+  call nonlocal ()
+
+contains
+
+  subroutine local ()
+    integer :: local_i
+    integer :: local_arg
+    integer :: local_a(N)
+    integer :: local_j
+
+    local_a (:) = 5
+    local_arg = 5
+
+    !$acc kernels loop &
+    !$acc gang(num:local_arg) worker(local_arg) vector(local_arg) &
+    !$acc wait async(local_arg)
+    do local_i = 1, N
+       !$acc cache (local_a(local_i:local_i + 5))
+       local_a(local_i) = 100
+       !$acc loop seq tile(*)
+       do local_j = 1, N
+       enddo
+       !$acc loop auto independent tile(1)
+       do local_j = 1, N
+       enddo
+    enddo
+    !$acc end kernels loop
+
+    !$acc kernels loop &
+    !$acc gang(static:local_arg) worker(local_arg) vector(local_arg) &
+    !$acc wait(local_arg, local_arg + 1, local_arg + 2) async
+    do local_i = 1, N
+       !$acc cache (local_a(local_i:local_i + 4))
+       local_a(local_i) = 100
+       !$acc loop seq tile(1)
+       do local_j = 1, N
+       enddo
+       !$acc loop auto independent tile(*)
+       do local_j = 1, N
+       enddo
+    enddo
+    !$acc end kernels loop
+  end subroutine local
+
+  subroutine nonlocal ()
+    nonlocal_a (:) = 5
+    nonlocal_arg = 5
+
+    !$acc kernels loop &
+    !$acc gang(num:nonlocal_arg) worker(nonlocal_arg) vector(nonlocal_arg) &
+    !$acc wait async(nonlocal_arg)
+    do nonlocal_i = 1, N
+       !$acc cache (nonlocal_a(nonlocal_i:nonlocal_i + 3))
+       nonlocal_a(nonlocal_i) = 100
+       !$acc loop seq tile(2)
+       do nonlocal_j = 1, N
+       enddo
+       !$acc loop auto independent tile(3)
+       do nonlocal_j = 1, N
+       enddo
+    enddo
+    !$acc end kernels loop
+
+    !$acc kernels loop &
+    !$acc gang(static:nonlocal_arg) worker(nonlocal_arg) vector(nonlocal_arg) &
+    !$acc wait(nonlocal_arg, nonlocal_arg + 1, nonlocal_arg + 2) async
+    do nonlocal_i = 1, N
+       !$acc cache (nonlocal_a(nonlocal_i:nonlocal_i + 2))
+       nonlocal_a(nonlocal_i) = 100
+       !$acc loop seq tile(*)
+       do nonlocal_j = 1, N
+       enddo
+       !$acc loop auto independent tile(*)
+       do nonlocal_j = 1, N
+       enddo
+    enddo
+    !$acc end kernels loop
+  end subroutine nonlocal
+end program main
diff --git a/gcc/testsuite/gfortran.dg/goacc/subroutines.f90 b/gcc/testsuite/gfortran.dg/goacc/subroutines.f90
deleted file mode 100644
index 6cab798d458..00000000000
--- a/gcc/testsuite/gfortran.dg/goacc/subroutines.f90
+++ /dev/null
@@ -1,73 +0,0 @@
-! Exercise how tree-nested.c handles gang, worker vector and seq.
-
-! { dg-do compile } 
-
-program main
-  integer, parameter :: N = 100
-  integer :: nonlocal_arg
-  integer :: nonlocal_a(N)
-  integer :: nonlocal_i
-  integer :: nonlocal_j
-  
-  nonlocal_a (:) = 5
-  nonlocal_arg = 5
-  
-  call local ()
-  call nonlocal ()
-
-contains
-
-  subroutine local ()
-    integer :: local_i
-    integer :: local_arg
-    integer :: local_a(N)
-    integer :: local_j
-    
-    local_a (:) = 5
-    local_arg = 5
-
-    !$acc kernels loop gang(num:local_arg) worker(local_arg) vector(local_arg)
-    do local_i = 1, N
-       local_a(local_i) = 100
-       !$acc loop seq
-       do local_j = 1, N
-       enddo
-    enddo
-    !$acc end kernels loop
-
-    !$acc kernels loop gang(static:local_arg) worker(local_arg) &
-    !$acc vector(local_arg)
-    do local_i = 1, N
-       local_a(local_i) = 100
-       !$acc loop seq
-       do local_j = 1, N
-       enddo
-    enddo
-    !$acc end kernels loop
-  end subroutine local
-
-  subroutine nonlocal ()
-    nonlocal_a (:) = 5
-    nonlocal_arg = 5
-  
-    !$acc kernels loop gang(num:nonlocal_arg) worker(nonlocal_arg) &
-    !$acc vector(nonlocal_arg)
-    do nonlocal_i = 1, N
-       nonlocal_a(nonlocal_i) = 100
-       !$acc loop seq
-       do nonlocal_j = 1, N
-       enddo
-    enddo
-    !$acc end kernels loop
-
-    !$acc kernels loop gang(static:nonlocal_arg) worker(nonlocal_arg) &
-    !$acc vector(nonlocal_arg)
-    do nonlocal_i = 1, N
-       nonlocal_a(nonlocal_i) = 100
-       !$acc loop seq
-       do nonlocal_j = 1, N
-       enddo
-    enddo
-    !$acc end kernels loop
-  end subroutine nonlocal
-end program main
diff --git a/gcc/tree-nested.c b/gcc/tree-nested.c
index 1433c3e2306..984fa818fa7 100644
--- a/gcc/tree-nested.c
+++ b/gcc/tree-nested.c
@@ -1114,6 +1114,8 @@ convert_nonlocal_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
 	case OMP_CLAUSE_GANG:
 	case OMP_CLAUSE_WORKER:
 	case OMP_CLAUSE_VECTOR:
+	case OMP_CLAUSE_ASYNC:
+	case OMP_CLAUSE_WAIT:
 	  /* Several OpenACC clauses have optional arguments.  Check if they
 	     are present.  */
 	  if (OMP_CLAUSE_OPERAND (clause, 0))
@@ -1197,9 +1199,27 @@ convert_nonlocal_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
 	case OMP_CLAUSE_SIMD:
 	case OMP_CLAUSE_DEFAULTMAP:
 	case OMP_CLAUSE_SEQ:
+	case OMP_CLAUSE_INDEPENDENT:
+	case OMP_CLAUSE_AUTO:
 	  break;
 
+	case OMP_CLAUSE_DEVICE_TYPE:
+	  /* TODO.  */
+	  gcc_unreachable ();
+
+	case OMP_CLAUSE_TILE:
+	  /* OpenACC tile clauses are discarded during gimplification, so we
+	     don't expect to see anything here.  */
+	  gcc_unreachable ();
+
+	case OMP_CLAUSE__CACHE_:
+	  /* These clauses belong to the OpenACC cache directive, which is
+	     discarded during gimplification, so we don't expect to see
+	     anything here.  */
+	  gcc_unreachable ();
+
 	case OMP_CLAUSE_BIND:
+	case OMP_CLAUSE_DEVICE_RESIDENT:
 	case OMP_CLAUSE_NOHOST:
 	default:
 	  gcc_unreachable ();
@@ -1792,6 +1812,8 @@ convert_local_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
 	case OMP_CLAUSE_GANG:
 	case OMP_CLAUSE_WORKER:
 	case OMP_CLAUSE_VECTOR:
+	case OMP_CLAUSE_ASYNC:
+	case OMP_CLAUSE_WAIT:
 	  /* Several OpenACC clauses have optional arguments.  Check if they
 	     are present.  */
 	  if (OMP_CLAUSE_OPERAND (clause, 0))
@@ -1880,9 +1902,27 @@ convert_local_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
 	case OMP_CLAUSE_SIMD:
 	case OMP_CLAUSE_DEFAULTMAP:
 	case OMP_CLAUSE_SEQ:
+	case OMP_CLAUSE_INDEPENDENT:
+	case OMP_CLAUSE_AUTO:
 	  break;
 
+	case OMP_CLAUSE_DEVICE_TYPE:
+	  /* TODO.  */
+	  gcc_unreachable ();
+
+	case OMP_CLAUSE_TILE:
+	  /* OpenACC tile clauses are discarded during gimplification, so we
+	     don't expect to see anything here.  */
+	  gcc_unreachable ();
+
+	case OMP_CLAUSE__CACHE_:
+	  /* These clauses belong to the OpenACC cache directive, which is
+	     discarded during gimplification, so we don't expect to see
+	     anything here.  */
+	  gcc_unreachable ();
+
 	case OMP_CLAUSE_BIND:
+	case OMP_CLAUSE_DEVICE_RESIDENT:
 	case OMP_CLAUSE_NOHOST:
 	default:
 	  gcc_unreachable ();