From 33718c02f44cf560c5725ed1681ae5981acbfc69 Mon Sep 17 00:00:00 2001
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Thu, 21 Mar 2019 20:13:44 +0000
Subject: [PATCH] [PR72741] Properly handle clauses specifying the level of
 parallelism for 'external' Fortran OpenACC routines

..., so as to also for these enable the generic middle end OMP code to verify
proper nesting of loops/routines regarding their levels of parallelism.

	gcc/fortran/
	PR fortran/72741
	* openmp.c (gfc_match_oacc_routine): Set the level of parallelism
	for all variants.
	(gfc_resolve_oacc_routines): Call gfc_add_omp_declare_target.
	gcc/testsuite/
	PR fortran/72741
	* c-c++-common/goacc/routine-3-extern.c: New file.
	* c-c++-common/goacc/routine-3.c: Adjust.
	* c-c++-common/goacc/routine-4-extern.c: New file.
	* c-c++-common/goacc/routine-4.c: Adjust.
	* gfortran.dg/goacc/routine-module-3.f90: New file.
	* gfortran.dg/goacc/routine-external-level-of-parallelism-1.f: New
	file.
	* gfortran.dg/goacc/routine-external-level-of-parallelism-2.f:
	Likewise.

git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@269858 138bc75d-0d04-0410-961f-82ee72b054a4
---
 gcc/fortran/ChangeLog                         |   5 +
 gcc/fortran/openmp.c                          |   8 +
 gcc/testsuite/ChangeLog                       |  16 +
 .../c-c++-common/goacc/routine-3-extern.c     |  89 +++++
 gcc/testsuite/c-c++-common/goacc/routine-3.c  |   1 +
 .../c-c++-common/goacc/routine-4-extern.c     | 124 ++++++
 gcc/testsuite/c-c++-common/goacc/routine-4.c  |   1 +
 .../routine-external-level-of-parallelism-1.f | 347 +++++++++++++++++
 .../routine-external-level-of-parallelism-2.f | 361 ++++++++++++++++++
 .../gfortran.dg/goacc/routine-module-3.f90    |  16 +
 10 files changed, 968 insertions(+)
 create mode 100644 gcc/testsuite/c-c++-common/goacc/routine-3-extern.c
 create mode 100644 gcc/testsuite/c-c++-common/goacc/routine-4-extern.c
 create mode 100644 gcc/testsuite/gfortran.dg/goacc/routine-external-level-of-parallelism-1.f
 create mode 100644 gcc/testsuite/gfortran.dg/goacc/routine-external-level-of-parallelism-2.f
 create mode 100644 gcc/testsuite/gfortran.dg/goacc/routine-module-3.f90

diff --git a/gcc/fortran/ChangeLog b/gcc/fortran/ChangeLog
index 7ce67eb46fe7..dd4347ef3d1b 100644
--- a/gcc/fortran/ChangeLog
+++ b/gcc/fortran/ChangeLog
@@ -1,5 +1,10 @@
 2019-03-21  Thomas Schwinge  <thomas@codesourcery.com>
 
+	PR fortran/72741
+	* openmp.c (gfc_match_oacc_routine): Set the level of parallelism
+	for all variants.
+	(gfc_resolve_oacc_routines): Call gfc_add_omp_declare_target.
+
 	PR fortran/89773
 	* gfortran.h (gfc_oacc_routine_name): Add loc member.
 	(gfc_resolve_oacc_routines): Declare.
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index 983b83db4a7b..9fc236760a1c 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -2391,6 +2391,8 @@ gfc_match_oacc_routine (void)
 
       if (add)
 	{
+	  sym->attr.oacc_routine_lop = lop;
+
 	  n = gfc_get_oacc_routine_name ();
 	  n->sym = sym;
 	  n->clauses = c;
@@ -6085,6 +6087,12 @@ gfc_resolve_oacc_routines (gfc_namespace *ns)
 		     " in !$ACC ROUTINE ( NAME ) at %L", sym->name, &orn->loc);
 	  continue;
 	}
+      if (!gfc_add_omp_declare_target (&sym->attr, sym->name, &orn->loc))
+	{
+	  gfc_error ("NAME %qs invalid"
+		     " in !$ACC ROUTINE ( NAME ) at %L", sym->name, &orn->loc);
+	  continue;
+	}
     }
 }
 
diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog
index e771a8743194..f575c0f59a9f 100644
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,4 +1,20 @@
 2019-03-21  Thomas Schwinge  <thomas@codesourcery.com>
+	    Cesar Philippidis  <cesar@codesourcery.com>
+
+	PR fortran/72741
+	* gfortran.dg/goacc/routine-external-level-of-parallelism-1.f: New
+	file.
+	* gfortran.dg/goacc/routine-external-level-of-parallelism-2.f:
+	Likewise.
+
+2019-03-21  Thomas Schwinge  <thomas@codesourcery.com>
+
+	PR fortran/72741
+	* c-c++-common/goacc/routine-3-extern.c: New file.
+	* c-c++-common/goacc/routine-3.c: Adjust.
+	* c-c++-common/goacc/routine-4-extern.c: New file.
+	* c-c++-common/goacc/routine-4.c: Adjust.
+	* gfortran.dg/goacc/routine-module-3.f90: New file.
 
 	PR fortran/89773
 	* gfortran.dg/goacc/pr89773.f90: New file.
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-3-extern.c b/gcc/testsuite/c-c++-common/goacc/routine-3-extern.c
new file mode 100644
index 000000000000..e32cfdefd2a2
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/routine-3-extern.c
@@ -0,0 +1,89 @@
+/* Test invalid calls to routines.  */
+/* Variant of 'routine-3.c', moving the callees 'extern'.  */
+
+#pragma acc routine gang
+extern int extern_gang (); /* { dg-message "declared here" "3" } */
+
+#pragma acc routine worker
+extern int extern_worker (); /* { dg-message "declared here" "2" } */
+
+#pragma acc routine vector
+extern int extern_vector (); /* { dg-message "declared here" } */
+
+#pragma acc routine seq
+extern int extern_seq ();
+
+int
+main ()
+{
+  int red = 0;
+#pragma acc parallel copy (red)
+  {
+    /* Independent/seq loop tests.  */
+#pragma acc loop reduction (+:red) // { dg-warning "insufficient partitioning" }
+    for (int i = 0; i < 10; i++)
+      red += extern_gang ();
+
+#pragma acc loop reduction (+:red)
+    for (int i = 0; i < 10; i++)
+      red += extern_worker ();
+
+#pragma acc loop reduction (+:red)
+    for (int i = 0; i < 10; i++)
+      red += extern_vector ();
+
+    /* Gang routine tests.  */
+#pragma acc loop gang reduction (+:red)  /* { dg-message "containing loop" } */
+    for (int i = 0; i < 10; i++)
+      red += extern_gang (); // { dg-error "routine call uses same" }
+
+#pragma acc loop worker reduction (+:red)  /* { dg-message "containing loop" } */
+    for (int i = 0; i < 10; i++)
+      red += extern_gang (); // { dg-error "routine call uses same" }
+
+#pragma acc loop vector reduction (+:red)  /* { dg-message "containing loop" } */
+    for (int i = 0; i < 10; i++)
+      red += extern_gang (); // { dg-error "routine call uses same" }
+
+    /* Worker routine tests.  */
+#pragma acc loop gang reduction (+:red)
+    for (int i = 0; i < 10; i++)
+      red += extern_worker ();
+
+#pragma acc loop worker reduction (+:red)  /* { dg-message "containing loop" } */
+    for (int i = 0; i < 10; i++)
+      red += extern_worker (); // { dg-error "routine call uses same" }
+
+#pragma acc loop vector reduction (+:red)  /* { dg-message "containing loop" } */
+    for (int i = 0; i < 10; i++)
+      red += extern_worker (); // { dg-error "routine call uses same" }
+
+    /* Vector routine tests.  */
+#pragma acc loop gang reduction (+:red)
+    for (int i = 0; i < 10; i++)
+      red += extern_vector ();
+
+#pragma acc loop worker reduction (+:red)
+    for (int i = 0; i < 10; i++)
+      red += extern_vector ();
+
+#pragma acc loop vector reduction (+:red)  /* { dg-message "containing loop" } */
+    for (int i = 0; i < 10; i++)
+      red += extern_vector (); // { dg-error "routine call uses same" }
+
+    /* Seq routine tests.  */
+#pragma acc loop gang reduction (+:red)
+    for (int i = 0; i < 10; i++)
+      red += extern_seq ();
+
+#pragma acc loop worker reduction (+:red)
+    for (int i = 0; i < 10; i++)
+      red += extern_seq ();
+
+#pragma acc loop vector reduction (+:red)
+    for (int i = 0; i < 10; i++)
+      red += extern_seq ();
+  }
+
+  return 0;
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-3.c b/gcc/testsuite/c-c++-common/goacc/routine-3.c
index eaea470fac09..364c8ad9ff5c 100644
--- a/gcc/testsuite/c-c++-common/goacc/routine-3.c
+++ b/gcc/testsuite/c-c++-common/goacc/routine-3.c
@@ -1,4 +1,5 @@
 /* Test invalid calls to routines.  */
+/* See also variant 'routine-3-extern.c', moving the callees 'extern'.  */
 
 #pragma acc routine gang
 int
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-4-extern.c b/gcc/testsuite/c-c++-common/goacc/routine-4-extern.c
new file mode 100644
index 000000000000..ec21db1c3194
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/routine-4-extern.c
@@ -0,0 +1,124 @@
+/* Test invalid intra-routine parallelism.  */
+/* Variant of 'routine-4.c', moving the callees 'extern'.  */
+
+extern void extern_gang (void);
+#pragma acc routine (extern_gang) gang
+extern void extern_worker (void);
+#pragma acc routine (extern_worker) worker
+extern void extern_vector (void);
+#pragma acc routine (extern_vector) vector
+extern void extern_seq (void);
+#pragma acc routine (extern_seq) seq
+
+void gang (void);
+void worker (void);
+void vector (void);
+
+#pragma acc routine (gang) gang
+#pragma acc routine (worker) worker
+#pragma acc routine (vector) vector
+  
+#pragma acc routine seq
+void seq (void)
+{
+  extern_gang ();  /* { dg-error "routine call uses" } */
+  extern_worker ();  /* { dg-error "routine call uses" } */
+  extern_vector ();  /* { dg-error "routine call uses" } */
+  extern_seq ();
+
+  int red;
+
+#pragma acc loop reduction (+:red) // { dg-warning "insufficient partitioning" }
+  for (int i = 0; i < 10; i++)
+    red ++;
+
+#pragma acc loop gang reduction (+:red) // { dg-error "disallowed by containing routine" }
+  for (int i = 0; i < 10; i++)
+    red ++;
+
+#pragma acc loop worker reduction (+:red) // { dg-error "disallowed by containing routine" }
+  for (int i = 0; i < 10; i++)
+    red ++;
+
+#pragma acc loop vector reduction (+:red) // { dg-error "disallowed by containing routine" }
+  for (int i = 0; i < 10; i++)
+    red ++;
+}
+
+void vector (void)
+{
+  extern_gang ();  /* { dg-error "routine call uses" } */
+  extern_worker ();  /* { dg-error "routine call uses" } */
+  extern_vector ();
+  extern_seq ();
+
+  int red;
+
+#pragma acc loop reduction (+:red)
+  for (int i = 0; i < 10; i++)
+    red ++;
+
+#pragma acc loop gang reduction (+:red) // { dg-error "disallowed by containing routine" }
+  for (int i = 0; i < 10; i++)
+    red ++;
+
+#pragma acc loop worker reduction (+:red) // { dg-error "disallowed by containing routine" }
+  for (int i = 0; i < 10; i++)
+    red ++;
+
+#pragma acc loop vector reduction (+:red)
+  for (int i = 0; i < 10; i++)
+    red ++;
+}
+
+void worker (void)
+{
+  extern_gang ();  /* { dg-error "routine call uses" } */
+  extern_worker ();
+  extern_vector ();
+  extern_seq ();
+
+  int red;
+
+#pragma acc loop reduction (+:red)
+  for (int i = 0; i < 10; i++)
+    red ++;
+
+#pragma acc loop gang reduction (+:red) // { dg-error "disallowed by containing routine" }
+  for (int i = 0; i < 10; i++)
+    red ++;
+
+#pragma acc loop worker reduction (+:red)
+  for (int i = 0; i < 10; i++)
+    red ++;
+
+#pragma acc loop vector reduction (+:red)
+  for (int i = 0; i < 10; i++)
+    red ++;
+}
+
+void gang (void)
+{
+  extern_gang ();
+  extern_worker ();
+  extern_vector ();
+  extern_seq ();
+
+  int red;
+
+#pragma acc loop reduction (+:red)
+  for (int i = 0; i < 10; i++)
+    red ++;
+
+#pragma acc loop gang reduction (+:red)
+  for (int i = 0; i < 10; i++)
+    red ++;
+
+#pragma acc loop worker reduction (+:red)
+  for (int i = 0; i < 10; i++)
+    red ++;
+
+#pragma acc loop vector reduction (+:red)
+  for (int i = 0; i < 10; i++)
+    red ++;
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-4.c b/gcc/testsuite/c-c++-common/goacc/routine-4.c
index efc4a0b95e59..5f2194c3f623 100644
--- a/gcc/testsuite/c-c++-common/goacc/routine-4.c
+++ b/gcc/testsuite/c-c++-common/goacc/routine-4.c
@@ -1,4 +1,5 @@
 /* Test invalid intra-routine parallelism.  */
+/* See also variant 'routine-4-extern.c', moving the callees 'extern'.  */
 
 void gang (void);
 void worker (void);
diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-external-level-of-parallelism-1.f b/gcc/testsuite/gfortran.dg/goacc/routine-external-level-of-parallelism-1.f
new file mode 100644
index 000000000000..c27fe7924ed5
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-external-level-of-parallelism-1.f
@@ -0,0 +1,347 @@
+! Check valid calls to 'external' OpenACC routines.
+
+! { dg-additional-options "-fopt-info-optimized-omp" }
+
+      subroutine sub
+      implicit none
+      integer, parameter :: n = 100
+      integer :: a(n), i, j
+      external :: gangr, workerr, vectorr, seqr
+!$acc routine (gangr) gang
+!$acc routine (workerr) worker
+!$acc routine (vectorr) vector
+!$acc routine (seqr) seq
+
+!
+! Test subroutine calls inside nested loops.
+!
+
+!$acc parallel loop ! { dg-message "optimized: assigned OpenACC gang loop parallelism" }
+      do i = 1, n
+!$acc loop ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
+         do j = 1, n
+            call workerr (a, n) ! { dg-message "optimized: assigned OpenACC worker vector loop parallelism" }
+         end do
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
+      do i = 1, n
+!$acc loop gang ! { dg-message "optimized: assigned OpenACC gang loop parallelism" }
+         do j = 1, n
+            call workerr (a, n) ! { dg-message "optimized: assigned OpenACC worker vector loop parallelism" }
+         end do
+      end do
+!$acc end parallel loop
+
+!
+! Test calls to seq routines
+!
+
+!$acc parallel loop ! { dg-message "optimized: assigned OpenACC gang vector loop parallelism" }
+      do i = 1, n
+         call seqr (a, n) ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop gang ! { dg-message "optimized: assigned OpenACC gang loop parallelism" }
+      do i = 1, n
+         call seqr (a, n) ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop worker ! { dg-message "optimized: assigned OpenACC worker loop parallelism" }
+      do i = 1, n
+         call seqr (a, n) ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop vector ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+      do i = 1, n
+         call seqr (a, n) ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop seq ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      do i = 1, n
+         call seqr (a, n) ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!
+! Test calls to gang routines
+!
+
+!$acc parallel loop ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
+      do i = 1, n
+         call gangr (a, n) ! { dg-message "optimized: assigned OpenACC gang worker vector loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop gang ! { dg-message "optimized: assigned OpenACC gang loop parallelism" }
+      do i = 1, n
+         call workerr (a, n) ! { dg-message "optimized: assigned OpenACC worker vector loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop worker ! { dg-message "optimized: assigned OpenACC worker loop parallelism" }
+      do i = 1, n
+         call seqr (a, n) ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop vector ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+      do i = 1, n
+         call seqr (a, n) ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop seq ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      do i = 1, n
+         call gangr (a, n) ! { dg-message "optimized: assigned OpenACC gang worker vector loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!
+! Test calls to worker routines
+!
+
+!$acc parallel loop ! { dg-message "optimized: assigned OpenACC gang loop parallelism" }
+      do i = 1, n
+         call workerr (a, n) ! { dg-message "optimized: assigned OpenACC worker vector loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop gang ! { dg-message "optimized: assigned OpenACC gang loop parallelism" }
+      do i = 1, n
+         call workerr (a, n) ! { dg-message "optimized: assigned OpenACC worker vector loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop worker ! { dg-message "optimized: assigned OpenACC worker loop parallelism" }
+      do i = 1, n
+         call vectorr (a, n) ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop vector ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+      do i = 1, n
+         call seqr (a, n) ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop seq ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      do i = 1, n
+         call workerr (a, n) ! { dg-message "optimized: assigned OpenACC worker vector loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!
+! Test calls to vector routines
+!
+
+!$acc parallel loop ! { dg-message "optimized: assigned OpenACC gang worker loop parallelism" }
+      do i = 1, n
+         call vectorr (a, n) ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop gang ! { dg-message "optimized: assigned OpenACC gang loop parallelism" }
+      do i = 1, n
+         call vectorr (a, n) ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop worker ! { dg-message "optimized: assigned OpenACC worker loop parallelism" }
+      do i = 1, n
+         call vectorr (a, n) ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop vector ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+      do i = 1, n
+         call seqr (a, n) ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop seq ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      do i = 1, n
+         call vectorr (a, n) ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+      end do
+!$acc end parallel loop
+      end subroutine sub
+
+      subroutine func
+      implicit none
+      integer, parameter :: n = 100
+      integer :: a(n), i, j
+      integer, external :: gangf, workerf, vectorf, seqf
+!$acc routine (gangf) gang
+!$acc routine (workerf) worker
+!$acc routine (vectorf) vector
+!$acc routine (seqf) seq
+
+!
+! Test subroutine calls inside nested loops.
+!
+
+!$acc parallel loop ! { dg-message "optimized: assigned OpenACC gang loop parallelism" }
+      do i = 1, n
+!$acc loop ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
+         do j = 1, n
+            a(i) = workerf (a, n) ! { dg-message "optimized: assigned OpenACC worker vector loop parallelism" }
+         end do
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
+      do i = 1, n
+!$acc loop gang ! { dg-message "optimized: assigned OpenACC gang loop parallelism" }
+         do j = 1, n
+            a(i) = vectorf (a, n) ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+         end do
+      end do
+!$acc end parallel loop
+
+!
+! Test calls to seq routines
+!
+
+!$acc parallel loop ! { dg-message "optimized: assigned OpenACC gang vector loop parallelism" }
+      do i = 1, n
+         a(i) = seqf (a, n) ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop gang ! { dg-message "optimized: assigned OpenACC gang loop parallelism" }
+      do i = 1, n
+         a(i) = seqf (a, n) ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop worker ! { dg-message "optimized: assigned OpenACC worker loop parallelism" }
+      do i = 1, n
+         a(i) = seqf (a, n) ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop vector ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+      do i = 1, n
+         a(i) = seqf (a, n) ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop seq ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      do i = 1, n
+         a(i) = seqf (a, n) ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!
+! Test calls to gang routines
+!
+
+!$acc parallel loop ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
+      do i = 1, n
+         a(i) = gangf (a, n) ! { dg-message "optimized: assigned OpenACC gang worker vector loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop gang ! { dg-message "optimized: assigned OpenACC gang loop parallelism" }
+      do i = 1, n
+         a(i) = seqf (a, n) ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop worker ! { dg-message "optimized: assigned OpenACC worker loop parallelism" }
+      do i = 1, n
+         a(i) = vectorf (a, n) ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop vector ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+      do i = 1, n
+         a(i) = seqf (a, n) ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop seq ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      do i = 1, n
+         a(i) = gangf (a, n) ! { dg-message "optimized: assigned OpenACC gang worker vector loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!
+! Test calls to worker routines
+!
+
+!$acc parallel loop ! { dg-message "optimized: assigned OpenACC gang loop parallelism" }
+      do i = 1, n
+         a(i) = workerf (a, n) ! { dg-message "optimized: assigned OpenACC worker vector loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop gang ! { dg-message "optimized: assigned OpenACC gang loop parallelism" }
+      do i = 1, n
+         a(i) = workerf (a, n) ! { dg-message "optimized: assigned OpenACC worker vector loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop worker ! { dg-message "optimized: assigned OpenACC worker loop parallelism" }
+      do i = 1, n
+         a(i) = vectorf (a, n) ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop vector ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+      do i = 1, n
+         a(i) = seqf (a, n) ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop seq ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      do i = 1, n
+         a(i) = workerf (a, n) ! { dg-message "optimized: assigned OpenACC worker vector loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!
+! Test calls to vector routines
+!
+
+!$acc parallel loop ! { dg-message "optimized: assigned OpenACC gang worker loop parallelism" }
+      do i = 1, n
+         a(i) = vectorf (a, n) ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop gang ! { dg-message "optimized: assigned OpenACC gang loop parallelism" }
+      do i = 1, n
+         a(i) = vectorf (a, n) ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop worker ! { dg-message "optimized: assigned OpenACC worker loop parallelism" }
+      do i = 1, n
+         a(i) = vectorf (a, n) ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop vector ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+      do i = 1, n
+         a(i) = seqf (a, n) ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop seq ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      do i = 1, n
+         a(i) = vectorf (a, n) ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+      end do
+!$acc end parallel loop
+      end subroutine func
diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-external-level-of-parallelism-2.f b/gcc/testsuite/gfortran.dg/goacc/routine-external-level-of-parallelism-2.f
new file mode 100644
index 000000000000..0e8dfb19e2b4
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-external-level-of-parallelism-2.f
@@ -0,0 +1,361 @@
+! Check invalid calls to 'external' OpenACC routines.
+
+! { dg-additional-options "-fopt-info-optimized-omp" }
+
+      subroutine sub
+      implicit none
+      integer, parameter :: n = 100
+      integer :: a(n), i, j
+      external :: gangr, workerr, vectorr, seqr
+!$acc routine (gangr) gang
+!$acc routine (workerr) worker
+!$acc routine (vectorr) vector
+!$acc routine (seqr) seq
+
+!
+! Test subroutine calls inside nested loops.
+!
+
+!$acc parallel loop ! { dg-message "optimized: assigned OpenACC gang loop parallelism" }
+      do i = 1, n
+!$acc loop ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
+         do j = 1, n
+            call workerr (a, n) ! { dg-message "optimized: assigned OpenACC worker vector loop parallelism" }
+         end do
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
+      do i = 1, n
+!$acc loop gang ! { dg-message "optimized: assigned OpenACC gang loop parallelism" }
+         do j = 1, n
+            call gangr (a, n) ! { dg-message "optimized: assigned OpenACC worker vector loop parallelism" }
+! { dg-error "routine call uses same OpenACC parallelism as containing loop" "" { target *-*-* } .-1 }
+         end do
+      end do
+!$acc end parallel loop
+
+!
+! Test calls to seq routines
+!
+
+!$acc parallel loop ! { dg-message "optimized: assigned OpenACC gang vector loop parallelism" }
+      do i = 1, n
+         call seqr (a, n) ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop gang ! { dg-message "optimized: assigned OpenACC gang loop parallelism" }
+      do i = 1, n
+         call seqr (a, n) ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop worker ! { dg-message "optimized: assigned OpenACC worker loop parallelism" }
+      do i = 1, n
+         call seqr (a, n) ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop vector ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+      do i = 1, n
+         call seqr (a, n) ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop seq ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      do i = 1, n
+         call seqr (a, n) ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!
+! Test calls to gang routines
+!
+
+!$acc parallel loop ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
+      do i = 1, n
+         call gangr (a, n) ! { dg-message "optimized: assigned OpenACC gang worker vector loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop gang ! { dg-message "optimized: assigned OpenACC gang loop parallelism" }
+      do i = 1, n
+         call gangr (a, n) ! { dg-message "optimized: assigned OpenACC worker vector loop parallelism" }
+! { dg-error "routine call uses same OpenACC parallelism as containing loop" "" { target *-*-* } .-1 }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop worker ! { dg-message "optimized: assigned OpenACC worker loop parallelism" }
+      do i = 1, n
+         call gangr (a, n) ! { dg-message "optimized: assigned OpenACC gang vector loop parallelism" }
+! { dg-error "routine call uses same OpenACC parallelism as containing loop" "" { target *-*-* } .-1 }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop vector ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+      do i = 1, n
+         call gangr (a, n) ! { dg-message "optimized: assigned OpenACC gang worker loop parallelism" }
+! { dg-error "routine call uses same OpenACC parallelism as containing loop" "" { target *-*-* } .-1 }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop seq ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      do i = 1, n
+         call gangr (a, n) ! { dg-message "optimized: assigned OpenACC gang worker vector loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!
+! Test calls to worker routines
+!
+
+!$acc parallel loop ! { dg-message "optimized: assigned OpenACC gang loop parallelism" }
+      do i = 1, n
+         call workerr (a, n) ! { dg-message "optimized: assigned OpenACC worker vector loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop gang ! { dg-message "optimized: assigned OpenACC gang loop parallelism" }
+      do i = 1, n
+         call workerr (a, n) ! { dg-message "optimized: assigned OpenACC worker vector loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop worker ! { dg-message "optimized: assigned OpenACC worker loop parallelism" }
+      do i = 1, n
+         call workerr (a, n) ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+! { dg-error "routine call uses same OpenACC parallelism as containing loop" "" { target *-*-* } .-1 }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop vector ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+      do i = 1, n
+         call workerr (a, n) ! { dg-message "optimized: assigned OpenACC worker loop parallelism" }
+! { dg-error "routine call uses same OpenACC parallelism as containing loop" "" { target *-*-* } .-1 }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop seq ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      do i = 1, n
+         call workerr (a, n) ! { dg-message "optimized: assigned OpenACC worker vector loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!
+! Test calls to vector routines
+!
+
+!$acc parallel loop ! { dg-message "optimized: assigned OpenACC gang worker loop parallelism" }
+      do i = 1, n
+         call vectorr (a, n) ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop gang ! { dg-message "optimized: assigned OpenACC gang loop parallelism" }
+      do i = 1, n
+         call vectorr (a, n) ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop worker ! { dg-message "optimized: assigned OpenACC worker loop parallelism" }
+      do i = 1, n
+         call vectorr (a, n) ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop vector ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+      do i = 1, n
+         call vectorr (a, n) ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+! { dg-error "routine call uses same OpenACC parallelism as containing loop" "" { target *-*-* } .-1 }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop seq ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      do i = 1, n
+         call vectorr (a, n) ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+      end do
+!$acc end parallel loop
+      end subroutine sub
+
+      subroutine func
+      implicit none
+      integer, parameter :: n = 100
+      integer :: a(n), i, j
+      integer, external :: gangf, workerf, vectorf, seqf
+!$acc routine (gangf) gang
+!$acc routine (workerf) worker
+!$acc routine (vectorf) vector
+!$acc routine (seqf) seq
+
+!
+! Test subroutine calls inside nested loops.
+!
+
+!$acc parallel loop ! { dg-message "optimized: assigned OpenACC gang loop parallelism" }
+      do i = 1, n
+!$acc loop ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
+         do j = 1, n
+            a(i) = workerf (a, n) ! { dg-message "optimized: assigned OpenACC worker vector loop parallelism" }
+         end do
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
+      do i = 1, n
+!$acc loop gang ! { dg-message "optimized: assigned OpenACC gang loop parallelism" }
+         do j = 1, n
+            a(i) = gangf (a, n) ! { dg-message "optimized: assigned OpenACC worker vector loop parallelism" }
+! { dg-error "routine call uses same OpenACC parallelism as containing loop" "" { target *-*-* } .-1 }
+         end do
+      end do
+!$acc end parallel loop
+
+!
+! Test calls to seq routines
+!
+
+!$acc parallel loop ! { dg-message "optimized: assigned OpenACC gang vector loop parallelism" }
+      do i = 1, n
+         a(i) = seqf (a, n) ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop gang ! { dg-message "optimized: assigned OpenACC gang loop parallelism" }
+      do i = 1, n
+         a(i) = seqf (a, n) ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop worker ! { dg-message "optimized: assigned OpenACC worker loop parallelism" }
+      do i = 1, n
+         a(i) = seqf (a, n) ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop vector ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+      do i = 1, n
+         a(i) = seqf (a, n) ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop seq ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      do i = 1, n
+         a(i) = seqf (a, n) ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!
+! Test calls to gang routines
+!
+
+!$acc parallel loop ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 }
+      do i = 1, n
+         a(i) = gangf (a, n) ! { dg-message "optimized: assigned OpenACC gang worker vector loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop gang ! { dg-message "optimized: assigned OpenACC gang loop parallelism" }
+      do i = 1, n
+         a(i) = gangf (a, n) ! { dg-message "optimized: assigned OpenACC worker vector loop parallelism" }
+! { dg-error "routine call uses same OpenACC parallelism as containing loop" "" { target *-*-* } .-1 }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop worker ! { dg-message "optimized: assigned OpenACC worker loop parallelism" }
+      do i = 1, n
+         a(i) = gangf (a, n) ! { dg-message "optimized: assigned OpenACC gang vector loop parallelism" }
+! { dg-error "routine call uses same OpenACC parallelism as containing loop" "" { target *-*-* } .-1 }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop vector ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+      do i = 1, n
+         a(i) = gangf (a, n) ! { dg-message "optimized: assigned OpenACC gang worker loop parallelism" }
+! { dg-error "routine call uses same OpenACC parallelism as containing loop" "" { target *-*-* } .-1 }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop seq ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      do i = 1, n
+         a(i) = gangf (a, n) ! { dg-message "optimized: assigned OpenACC gang worker vector loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!
+! Test calls to worker routines
+!
+
+!$acc parallel loop ! { dg-message "optimized: assigned OpenACC gang loop parallelism" }
+      do i = 1, n
+         a(i) = workerf (a, n) ! { dg-message "optimized: assigned OpenACC worker vector loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop gang ! { dg-message "optimized: assigned OpenACC gang loop parallelism" }
+      do i = 1, n
+         a(i) = workerf (a, n) ! { dg-message "optimized: assigned OpenACC worker vector loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop worker ! { dg-message "optimized: assigned OpenACC worker loop parallelism" }
+      do i = 1, n
+         a(i) = workerf (a, n) ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+! { dg-error "routine call uses same OpenACC parallelism as containing loop" "" { target *-*-* } .-1 }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop vector ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+      do i = 1, n
+         a(i) = workerf (a, n) ! { dg-message "optimized: assigned OpenACC worker loop parallelism" }
+! { dg-error "routine call uses same OpenACC parallelism as containing loop" "" { target *-*-* } .-1 }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop seq ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      do i = 1, n
+         a(i) = workerf (a, n) ! { dg-message "optimized: assigned OpenACC worker vector loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!
+! Test calls to vector routines
+!
+
+!$acc parallel loop ! { dg-message "optimized: assigned OpenACC gang worker loop parallelism" }
+      do i = 1, n
+         a(i) = vectorf (a, n) ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop gang ! { dg-message "optimized: assigned OpenACC gang loop parallelism" }
+      do i = 1, n
+         a(i) = vectorf (a, n) ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop worker ! { dg-message "optimized: assigned OpenACC worker loop parallelism" }
+      do i = 1, n
+         a(i) = vectorf (a, n) ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop vector ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+      do i = 1, n
+         a(i) = vectorf (a, n) ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+! { dg-error "routine call uses same OpenACC parallelism as containing loop" "" { target *-*-* } .-1 }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop seq ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
+      do i = 1, n
+         a(i) = vectorf (a, n) ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
+      end do
+!$acc end parallel loop
+      end subroutine func
diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-module-3.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-module-3.f90
new file mode 100644
index 000000000000..a4ff54954afa
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-module-3.f90
@@ -0,0 +1,16 @@
+! Invalid use of routines defined inside a Fortran module.
+
+! { dg-compile-aux-modules "routine-module-mod-1.f90" }
+
+program main
+  use routine_module_mod_1
+  implicit none
+  !$acc routine (s_1) seq ! { dg-error "Cannot change attributes of USE-associated symbol s_1" }
+   ! { dg-error "NAME 's_1' invalid in \\!\\\$ACC ROUTINE \\( NAME \\)" "" { target *-*-* } .-1 }
+  !$acc routine (s_2) seq ! { dg-error "Cannot change attributes of USE-associated symbol s_2" }
+   ! { dg-error "NAME 's_2' invalid in \\!\\\$ACC ROUTINE \\( NAME \\)" "" { target *-*-* } .-1 }
+  !$acc routine (v_1) seq ! { dg-error "Cannot change attributes of USE-associated symbol v_1" }
+   ! { dg-error "NAME 'v_1' invalid in \\!\\\$ACC ROUTINE \\( NAME \\)" "" { target *-*-* } .-1 }
+  !$acc routine (w_1) gang ! { dg-error "Cannot change attributes of USE-associated symbol w_1" }
+   ! { dg-error "NAME 'w_1' invalid in \\!\\\$ACC ROUTINE \\( NAME \\)" "" { target *-*-* } .-1 }
+end program main
-- 
2.17.1

