[PR72741] Properly handle clauses specifying the level of parallelism for…

[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.

Co-Authored-By: Cesar Philippidis <cesar@codesourcery.com>

From-SVN: r269858

gcc/fortran/ChangeLog

 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.

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;
+	}
     }
 }
 

gcc/testsuite/ChangeLog

 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.

gcc/testsuite/c-c++-common/goacc/routine-3-extern.c

+/* 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;
+}

gcc/testsuite/c-c++-common/goacc/routine-3.c

 /* Test invalid calls to routines.  */
+/* See also variant 'routine-3-extern.c', moving the callees 'extern'.  */
 
 #pragma acc routine gang
 int

gcc/testsuite/c-c++-common/goacc/routine-4-extern.c

+/* 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 ++;
+}

gcc/testsuite/c-c++-common/goacc/routine-4.c

 /* Test invalid intra-routine parallelism.  */
+/* See also variant 'routine-4-extern.c', moving the callees 'extern'.  */
 
 void gang (void);
 void worker (void);

gcc/testsuite/gfortran.dg/goacc/routine-external-level-of-parallelism-1.f

+! 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

gcc/testsuite/gfortran.dg/goacc/routine-external-level-of-parallelism-2.f

+! 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

gcc/testsuite/gfortran.dg/goacc/routine-module-3.f90

+! 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