diff mbox

OpenACC routines -- test cases

Message ID 6b840a02-5355-5e87-98e8-21968429589c@codesourcery.com
State New
Headers show

Commit Message

Cesar Philippidis Nov. 11, 2016, 11:44 p.m. UTC
This patch contains new and adjusted, runtime and compiler test cases
for the new OpenACC routine functionality.

Is this ok for trunk?

Cesar

Comments

Jakub Jelinek Nov. 18, 2016, 12:30 p.m. UTC | #1
On Fri, Nov 11, 2016 at 03:44:29PM -0800, Cesar Philippidis wrote:
> This patch contains new and adjusted, runtime and compiler test cases
> for the new OpenACC routine functionality.
> 
> Is this ok for trunk?

Ok (though of course, if the diagnostic wording is adjusted, then
the test will need to match).

	Jakub
diff mbox

Patch

2016-11-11  Cesar Philippidis  <cesar@codesourcery.com>
	    Thomas Schwinge  <thomas@codesourcery.com>

	gcc/testsuite/
	* c-c++-common/goacc/routine-1.c: Add more coverage.
	* c-c++-common/goacc/routine-2.c: Likewise.
	* c-c++-common/goacc/routine-5.c: Likewise.
	* c-c++-common/goacc/routine-level-of-parallelism-1.c: New test.
	* c-c++-common/goacc/routine-level-of-parallelism-2.c: New test.
	* c-c++-common/goacc/routine-nohost-1.c: New test.
	* c-c++-common/goacc/routine-nohost-2.c: New test.
	* gfortran.dg/goacc/fixed-1.f: Add more coverage.
	* gfortran.dg/goacc/loop-5.f95: Likewise.
	* gfortran.dg/goacc/routine-6.f90: Likewise.
	* gfortran.dg/goacc/routine-7.f90: New test.
	* gfortran.dg/goacc/routine-9.f90: New test.
	* gfortran.dg/goacc/routine-nested-parallelism.f: New test.
	* gfortran.dg/goacc/routine-nested-parallelism.f90: New test.

	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/routine-3.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/routine-5.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/routine-bind-nohost-1.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/routine-nohost-1.c: New test.
	* testsuite/libgomp.oacc-fortran/abort-1.f90: Add more coverage.
	* testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f: Likewise.
	* testsuite/libgomp.oacc-fortran/routine-6.f90: New test.
	* testsuite/libgomp.oacc-fortran/routine-7.f90: Add more coverage.
	* testsuite/libgomp.oacc-fortran/routine-8.f90: New test.
	* testsuite/libgomp.oacc-fortran/vector-routine.f90: New test.


diff --git a/gcc/testsuite/c-c++-common/goacc/routine-1.c b/gcc/testsuite/c-c++-common/goacc/routine-1.c
index a5e0d69..43434a3 100644
--- a/gcc/testsuite/c-c++-common/goacc/routine-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/routine-1.c
@@ -1,3 +1,4 @@ 
+/* Test valid use of clauses with routine.  */
 
 #pragma acc routine gang
 void gang (void)
@@ -19,15 +20,88 @@  void seq (void)
 {
 }
 
-int main ()
+
+#pragma acc routine
+void bind_f_1 (void)
+{
+}
+
+#pragma acc routine
+extern void bind_f_1 (void);
+
+#pragma acc routine (bind_f_1)
+
+
+#pragma acc routine \
+  bind (bind_f_1)
+void bind_f_1_1 (void)
+{
+}
+
+#pragma acc routine \
+  bind (bind_f_1)
+extern void bind_f_1_1 (void);
+
+#pragma acc routine (bind_f_1_1) \
+  bind (bind_f_1)
+
+
+/* Non-sensical bind clause, but permitted.  */
+#pragma acc routine \
+  bind ("bind_f_2")
+void bind_f_2 (void)
+{
+}
+
+#pragma acc routine \
+  bind ("bind_f_2")
+extern void bind_f_2 (void);
+
+#pragma acc routine (bind_f_2) \
+  bind ("bind_f_2")
+
+
+#pragma acc routine \
+  bind ("bind_f_2")
+void bind_f_2_1 (void)
+{
+}
+
+#pragma acc routine \
+  bind ("bind_f_2")
+extern void bind_f_2_1 (void);
+
+#pragma acc routine (bind_f_2_1) \
+  bind ("bind_f_2")
+
+
+#pragma acc routine \
+  nohost
+void nohost (void)
 {
+}
+
+#pragma acc routine \
+  nohost
+extern void nohost (void);
 
+#pragma acc routine (nohost) \
+  nohost
+
+
+int main ()
+{
 #pragma acc parallel num_gangs (32) num_workers (32) vector_length (32)
   {
     gang ();
     worker ();
     vector ();
     seq ();
+    bind_f_1 ();
+    bind_f_1_1 ();
+    bind_f_2 ();
+    bind_f_2_1 ();
+    nohost ();
   }
 
   return 0;
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-2.c b/gcc/testsuite/c-c++-common/goacc/routine-2.c
index fc5eb11..debf6d7 100644
--- a/gcc/testsuite/c-c++-common/goacc/routine-2.c
+++ b/gcc/testsuite/c-c++-common/goacc/routine-2.c
@@ -1,21 +1,171 @@ 
-#pragma acc routine gang worker /* { dg-error "multiple loop axes" } */
-void gang (void)
+/* Test invalid use of clauses with OpenACC routine.  */
+
+extern float F;
+#pragma acc routine bind (F) /* { dg-error ".F. does not refer to a function" } */
+extern void F_1 (void);
+
+typedef int T;
+#pragma acc routine bind (T) /* { dg-error ".T. does not refer to a function" } */
+extern void T_1 (void);
+
+#pragma acc routine (nothing) gang /* { dg-error ".nothing. has not been declared" } */
+
+#pragma acc routine bind (bind_0) /* { dg-error ".bind_0. has not been declared" }*/
+extern void bind_0 (void);
+
+extern void a(void), b(void);
+
+#pragma acc routine bind(a) bind(b) /* { dg-error "too many .bind. clauses" } */
+extern void bind_1 (void);
+
+#pragma acc routine bind(a) bind("b") /* { dg-error "too many .bind. clauses" } */
+extern void bind_2 (void);
+
+#pragma acc routine bind("a") bind("b") /* { dg-error "too many .bind. clauses" } */
+extern void bind_3 (void);
+
+#pragma acc routine nohost nohost /* { dg-error "too many .nohost. clauses" } */
+extern void nohost (void);
+
+
+/* bind clause on first OpenACC routine directive but not on following.  */
+
+extern void a_bind_f_1 (void);
+#pragma acc routine (a_bind_f_1)
+
+
+#pragma acc routine \
+  bind (a_bind_f_1)
+void a_bind_f_1_1 (void)
+{
+}
+
+#pragma acc routine /* { dg-error "missing .bind. clause when applying .#pragma acc routine. to .\[void \]*a_bind_f_1_1\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
+extern void a_bind_f_1_1 (void);
+
+#pragma acc routine (a_bind_f_1_1) /* { dg-error "missing .bind. clause when applying .#pragma acc routine. to .\[void \]*a_bind_f_1_1\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
+
+
+/* Non-sensical bind clause, but permitted.  */
+#pragma acc routine \
+  bind ("a_bind_f_2")
+void a_bind_f_2 (void)
 {
 }
 
-#pragma acc routine worker vector /* { dg-error "multiple loop axes" } */
-void worker (void)
+#pragma acc routine /* { dg-error "missing .bind. clause when applying .#pragma acc routine. to .\[void \]*a_bind_f_2\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
+extern void a_bind_f_2 (void);
+
+#pragma acc routine (a_bind_f_2) /* { dg-error "missing .bind. clause when applying .#pragma acc routine. to .\[void \]*a_bind_f_2\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
+
+
+#pragma acc routine \
+  bind ("a_bind_f_2")
+void a_bind_f_2_1 (void)
 {
 }
 
-#pragma acc routine vector seq /* { dg-error "multiple loop axes" } */
-void vector (void)
+#pragma acc routine /* { dg-error "missing .bind. clause when applying .#pragma acc routine. to .\[void \]*a_bind_f_2_1\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
+extern void a_bind_f_2_1 (void);
+
+#pragma acc routine (a_bind_f_2_1) /* { dg-error "missing .bind. clause when applying .#pragma acc routine. to .\[void \]*a_bind_f_2_1\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
+
+
+/* No bind clause on first OpenACC routine directive, but on following.  */
+
+#pragma acc routine
+extern void b_bind_f_1 (void);
+
+
+#pragma acc routine
+void b_bind_f_1_1 (void)
 {
 }
 
-#pragma acc routine seq gang /* { dg-error "multiple loop axes" } */
-void seq (void)
+#pragma acc routine \
+  bind (b_bind_f_1) /* { dg-error "incompatible .bind. clause when applying .#pragma acc routine. to .\[void \]*b_bind_f_1_1\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
+extern void b_bind_f_1_1 (void);
+
+#pragma acc routine (b_bind_f_1_1) \
+  bind (b_bind_f_1) /* { dg-error "incompatible .bind. clause when applying .#pragma acc routine. to .\[void \]*b_bind_f_1_1\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
+
+
+/* Non-sensical bind clause, but permitted.  */
+#pragma acc routine
+void b_bind_f_2 (void)
 {
 }
 
-#pragma acc routine (nothing) gang /* { dg-error "not been declared" } */
+#pragma acc routine \
+  bind ("b_bind_f_2") /* { dg-error "incompatible .bind. clause when applying .#pragma acc routine. to .\[void \]*b_bind_f_2\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
+extern void b_bind_f_2 (void);
+
+#pragma acc routine (b_bind_f_2) \
+  bind ("b_bind_f_2") /* { dg-error "incompatible .bind. clause when applying .#pragma acc routine. to .\[void \]*b_bind_f_2\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
+
+
+#pragma acc routine
+void b_bind_f_2_1 (void)
+{
+}
+
+#pragma acc routine \
+  bind ("b_bind_f_2") /* { dg-error "incompatible .bind. clause when applying .#pragma acc routine. to .\[void \]*b_bind_f_2_1\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
+extern void b_bind_f_2_1 (void);
+
+#pragma acc routine (b_bind_f_2_1) \
+  bind ("b_bind_f_2") /* { dg-error "incompatible .bind. clause when applying .#pragma acc routine. to .\[void \]*b_bind_f_2_1\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
+
+
+/* Non-matching bind clauses.  */
+
+#pragma acc routine
+void c_bind_f_1a (void)
+{
+}
+
+#pragma acc routine
+extern void c_bind_f_1b (void);
+
+
+#pragma acc routine \
+  bind (c_bind_f_1a)
+void c_bind_f_1_1 (void)
+{
+}
+
+#pragma acc routine \
+  bind (c_bind_f_1b) /* { dg-error "incompatible .bind. clause when applying .#pragma acc routine. to .\[void \]*c_bind_f_1_1\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
+extern void c_bind_f_1_1 (void);
+
+#pragma acc routine (c_bind_f_1_1) \
+  bind (c_bind_f_1b) /* { dg-error "incompatible .bind. clause when applying .#pragma acc routine. to .\[void \]*c_bind_f_1_1\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
+
+
+/* Non-sensical bind clause, but permitted.  */
+#pragma acc routine \
+  bind ("c_bind_f_2")
+void c_bind_f_2 (void)
+{
+}
+
+#pragma acc routine \
+  bind ("C_BIND_F_2") /* { dg-error "incompatible .bind. clause when applying .#pragma acc routine. to .\[void \]*c_bind_f_2\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
+extern void c_bind_f_2 (void);
+
+#pragma acc routine (c_bind_f_2) \
+  bind ("C_BIND_F_2") /* { dg-error "incompatible .bind. clause when applying .#pragma acc routine. to .\[void \]*c_bind_f_2\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
+
+
+#pragma acc routine \
+  bind ("c_bind_f_2")
+void c_bind_f_2_1 (void)
+{
+}
+
+#pragma acc routine \
+  bind ("C_BIND_F_2") /* { dg-error "incompatible .bind. clause when applying .#pragma acc routine. to .\[void \]*c_bind_f_2_1\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
+extern void c_bind_f_2_1 (void);
+
+#pragma acc routine (c_bind_f_2_1) \
+  bind ("C_BIND_F_2") /* { dg-error "incompatible .bind. clause when applying .#pragma acc routine. to .\[void \]*c_bind_f_2_1\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-5.c b/gcc/testsuite/c-c++-common/goacc/routine-5.c
index 17fe67c..f10651d 100644
--- a/gcc/testsuite/c-c++-common/goacc/routine-5.c
+++ b/gcc/testsuite/c-c++-common/goacc/routine-5.c
@@ -150,61 +150,19 @@  void f_static_assert();
 
 #pragma acc routine
 __extension__ extern void ex1();
-#pragma acc routine (ex1) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*ex1" } */
+#pragma acc routine (ex1) worker /* { dg-error "has already been marked as an accelerator routine" } */
 
 #pragma acc routine
 __extension__ __extension__ __extension__ __extension__ __extension__ void ex2()
 {
 }
-#pragma acc routine (ex2) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*ex2" } */
+#pragma acc routine (ex2) worker /* { dg-error "has already been marked as an accelerator routine" } */
 
 #pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
 __extension__ int ex3;
 #pragma acc routine (ex3) /* { dg-error ".ex3. does not refer to a function" } */
 
 
-/* "#pragma acc routine" already applied.  */
-
-extern void fungsi_1();
-#pragma acc routine(fungsi_1) gang
-#pragma acc routine(fungsi_1) gang /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_1" } */
-#pragma acc routine(fungsi_1) worker /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_1" } */
-#pragma acc routine(fungsi_1) vector /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_1" } */
-
-#pragma acc routine seq
-extern void fungsi_2();
-#pragma acc routine(fungsi_2) seq /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_2." } */
-#pragma acc routine(fungsi_2) worker /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_2." } */
-#pragma acc routine(fungsi_2) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_2." } */
-
-#pragma acc routine vector
-extern void fungsi_3();
-#pragma acc routine vector /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_3." } */
-void fungsi_3()
-{
-}
-
-extern void fungsi_4();
-#pragma acc routine (fungsi_4) worker
-#pragma acc routine gang /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_4." } */
-void fungsi_4()
-{
-}
-
-#pragma acc routine gang
-void fungsi_5()
-{
-}
-#pragma acc routine (fungsi_5) worker /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_5." } */
-
-#pragma acc routine seq
-void fungsi_6()
-{
-}
-#pragma acc routine seq /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_6." } */
-extern void fungsi_6();
-
-
 /* "#pragma acc routine" must be applied before.  */
 
 void Bar ();
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-1.c b/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-1.c
new file mode 100644
index 0000000..8f45499
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-1.c
@@ -0,0 +1,450 @@ 
+/* Test various aspects of clauses specifying incompatible levels of
+   parallelism with the OpenACC routine directive.  The Fortran counterpart is
+   ../../gfortran.dg/goacc/routine-level-of-parallelism-1.f.  */
+
+extern void g_1 (void);
+#pragma acc routine (g_1) gang gang /* { dg-error "too many 'gang' clauses" } */
+
+#pragma acc routine worker worker /* { dg-error "too many 'worker' clauses" } */
+void w_1 (void)
+{
+}
+
+#pragma acc routine vector vector /* { dg-error "too many 'vector' clauses" } */
+void v_1 (void)
+{
+}
+
+#pragma acc routine seq seq /* { dg-error "too many 'seq' clauses" } */
+extern void s_1 (void);
+
+
+#pragma acc routine gang gang gang /* { dg-error "too many 'gang' clauses" } */
+void g_2 (void)
+{
+}
+
+#pragma acc routine worker worker worker /* { dg-error "too many 'worker' clauses" } */
+extern void w_2 (void);
+
+extern void v_2 (void);
+#pragma acc routine (v_2) vector vector vector /* { dg-error "too many 'vector' clauses" } */
+
+#pragma acc routine seq seq seq /* { dg-error "too many 'seq' clauses" } */
+void s_2 (void)
+{
+}
+
+
+#pragma acc routine \
+  gang \
+  worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */
+void g_3 (void)
+{
+}
+#pragma acc routine (g_3) \
+  gang \
+  seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */
+#pragma acc routine (g_3) \
+  gang \
+  vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */
+
+extern void w_3 (void);
+#pragma acc routine (w_3) \
+  worker \
+  vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */
+#pragma acc routine (w_3) \
+  worker \
+  gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */
+#pragma acc routine (w_3) \
+  worker \
+  seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */
+
+#pragma acc routine \
+  vector \
+  seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */
+void v_3 (void)
+{
+}
+#pragma acc routine (v_3) \
+  vector \
+  worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */
+#pragma acc routine (v_3) \
+  vector \
+  gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */
+
+extern void s_3 (void);
+#pragma acc routine (s_3) \
+  seq \
+  gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */
+#pragma acc routine (s_3) \
+  seq \
+  vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */
+#pragma acc routine (s_3) \
+  seq \
+  worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */
+
+
+#pragma acc routine \
+  gang gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */ \
+  vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ \
+  seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */
+extern void g_4 (void);
+#pragma acc routine (g_4) \
+  gang gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ \
+  worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */ \
+  seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */
+#pragma acc routine (g_4) \
+  gang gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */ \
+  vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ \
+  worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */
+
+extern void w_4 (void);
+#pragma acc routine (w_4) \
+  worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */ \
+  vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ \
+  gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */
+#pragma acc routine (w_4) \
+  worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ \
+  seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */ \
+  gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */
+#pragma acc routine (w_4) \
+  worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */ \
+  gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */ \
+  vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */
+
+#pragma acc routine \
+  vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */ \
+  seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */ \
+  gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */
+void v_4 (void)
+{
+}
+#pragma acc routine (v_4) \
+  vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */ \
+  worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */ \
+  gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */
+#pragma acc routine (v_4) \
+  vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */ \
+  seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */ \
+  worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */
+
+#pragma acc routine \
+  seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ \
+  worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */ \
+  gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */
+void s_4 (void)
+{
+}
+#pragma acc routine (s_4) \
+  seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ \
+  gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */ \
+  worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */
+#pragma acc routine (s_4) \
+  seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */ \
+  vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ \
+  gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */
+
+
+#pragma acc routine \
+  gang gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } 163 } */ \
+  vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } 165 } */ \
+  seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } 167 } */
+void g_5 (void)
+{
+}
+#pragma acc routine (g_5) \
+  gang gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } 174 } */ \
+  worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } 176 } */ \
+  seq seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } 178 } */
+#pragma acc routine (g_5) \
+  gang gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } 182 } */ \
+  vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } 184 } */ \
+  worker worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } 186 } */
+
+#pragma acc routine \
+  worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } 191 } */ \
+  vector vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } 193 } */ \
+  gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } 195 } */
+extern void w_5 (void);
+#pragma acc routine (w_5) \
+  worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } 200 } */ \
+  seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } 202 } */ \
+  gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } 204 } */
+#pragma acc routine (w_5) \
+  worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } 208 } */ \
+  gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } 210 } */ \
+  vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } 212 } */
+
+#pragma acc routine \
+  vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } 217 } */ \
+  seq seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } 219 } */ \
+  gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } 221 } */
+extern void v_5 (void);
+#pragma acc routine (v_5) \
+  vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  seq seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } 226 } */ \
+  worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } 228 } */ \
+  gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } 230 } */
+#pragma acc routine (v_5) \
+  vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } 234 } */ \
+  seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } 236 } */ \
+  worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } 238 } */
+
+extern void s_5 (void);
+#pragma acc routine (s_5) \
+  seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } 244 } */ \
+  worker worker worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } 246 } */ \
+  gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } 248 } */
+#pragma acc routine (s_5) \
+  seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } 252 } */ \
+  gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } 254 } */ \
+  worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } 256 } */
+#pragma acc routine (s_5) \
+  seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } 260 } */ \
+  vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } 262 } */ \
+  gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } 264 } */
+
+
+/* Like the *_5 tests, but with the order of clauses changed in the second and
+   following routine directives for the specific *_5 function.  */
+
+#pragma acc routine \
+  gang gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } 273 } */ \
+  vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } 275 } */ \
+  seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } 277 } */
+void g_6 (void)
+{
+}
+#pragma acc routine (g_6) \
+  vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  /* { dg-error "incompatible .vector. clause when applying .#pragma acc routine. to .\[void \]*g_6\[\\(\\)\]*., which has already been marked as an accelerator routine" "" { target *-*-* } 283 } */ \
+  gang gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } 285 } */ \
+  worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } 287 } */ \
+  seq seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } 289 } */
+#pragma acc routine (g_6) \
+  seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  /* { dg-error "incompatible .seq. clause when applying .#pragma acc routine. to .\[void \]*g_6\[\\(\\)\]*., which has already been marked as an accelerator routine" "" { target *-*-* } 292 } */ \
+  gang gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } 294 } */ \
+  vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } 296 } */ \
+  worker worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } 298 } */
+
+#pragma acc routine \
+  worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } 303 } */ \
+  vector vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } 305 } */ \
+  gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } 307 } */
+extern void w_6 (void);
+#pragma acc routine (w_6) \
+  vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  /* { dg-error "incompatible .vector. clause when applying .#pragma acc routine. to .\[void \]*w_6\[\\(\\)\]*., which has already been marked as an accelerator routine" "" { target *-*-* } 311 } */ \
+  worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } 313 } */ \
+  seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } 315 } */ \
+  gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } 317 } */
+#pragma acc routine (w_6) \
+  seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  /* { dg-error "incompatible .seq. clause when applying .#pragma acc routine. to .\[void \]*w_6\[\\(\\)\]*., which has already been marked as an accelerator routine" "" { target *-*-* } 320 } */ \
+  worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } 322 } */ \
+  gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } 324 } */ \
+  vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } 326 } */
+
+#pragma acc routine \
+  vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } 331 } */ \
+  seq seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } 333 } */ \
+  gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } 335 } */
+extern void v_6 (void);
+#pragma acc routine (v_6) \
+  seq seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  /* { dg-error "incompatible .seq. clause when applying .#pragma acc routine. to .\[void \]*v_6\[\\(\\)\]*., which has already been marked as an accelerator routine" "" { target *-*-* } 339 } */ \
+  vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } 341 } */ \
+  worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } 343 } */ \
+  gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } 345 } */
+#pragma acc routine (v_6) \
+  gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  /* { dg-error "incompatible .gang. clause when applying .#pragma acc routine. to .\[void \]*v_6\[\\(\\)\]*., which has already been marked as an accelerator routine" "" { target *-*-* } 348 } */ \
+  vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } 350 } */ \
+  seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } 352 } */ \
+  worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } 354 } */
+
+extern void s_6 (void);
+#pragma acc routine (s_6) \
+  seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } 360 } */ \
+  worker worker worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } 362 } */ \
+  gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } 364 } */
+#pragma acc routine (s_6) \
+  vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  /* { dg-error "incompatible .vector. clause when applying .#pragma acc routine. to .\[void \]*s_6\[\\(\\)\]*., which has already been marked as an accelerator routine" "" { target *-*-* } 367 } */ \
+  seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } 369 } */ \
+  gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } 371 } */ \
+  worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } 373 } */
+#pragma acc routine (s_6) \
+  worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  /* { dg-error "incompatible .worker. clause when applying .#pragma acc routine. to .\[void \]*s_6\[\\(\\)\]*., which has already been marked as an accelerator routine" "" { target *-*-* } 376 } */ \
+  seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } 378 } */ \
+  vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } 380 } */ \
+  gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } 382 } */
+
+
+/* Like the *_6 tests, but without all the duplicate clauses, so that the
+   routine directives are valid in isolation.  */
+
+#pragma acc routine \
+  gang
+void g_7 (void)
+{
+}
+#pragma acc routine (g_7) \
+  vector /* { dg-error "incompatible .vector. clause when applying .#pragma acc routine. to .\[void \]*g_7\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
+#pragma acc routine (g_7) \
+  seq /* { dg-error "incompatible .seq. clause when applying .#pragma acc routine. to .\[void \]*g_7\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
+
+#pragma acc routine \
+  worker
+extern void w_7 (void);
+#pragma acc routine (w_7) \
+  vector /* { dg-error "incompatible .vector. clause when applying .#pragma acc routine. to .\[void \]*w_7\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
+#pragma acc routine (w_7) \
+  seq /* { dg-error "incompatible .seq. clause when applying .#pragma acc routine. to .\[void \]*w_7\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
+
+#pragma acc routine \
+  vector
+extern void v_7 (void);
+#pragma acc routine (v_7) \
+  seq /* { dg-error "incompatible .seq. clause when applying .#pragma acc routine. to .\[void \]*v_7\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
+#pragma acc routine (v_7) \
+  gang /* { dg-error "incompatible .gang. clause when applying .#pragma acc routine. to .\[void \]*v_7\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
+
+extern void s_7 (void);
+#pragma acc routine (s_7) \
+  seq
+#pragma acc routine (s_7) \
+  vector /* { dg-error "incompatible .vector. clause when applying .#pragma acc routine. to .\[void \]*s_7\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
+#pragma acc routine (s_7) \
+  worker /* { dg-error "incompatible .worker. clause when applying .#pragma acc routine. to .\[void \]*s_7\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
+
+
+/* Test cases for implicit seq clause.  */
+
+#pragma acc routine \
+  gang
+void g_8 (void)
+{
+}
+#pragma acc routine (g_8) /* { dg-error "incompatible .seq. clause when applying .#pragma acc routine. to .\[void \]*g_8\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
+
+#pragma acc routine \
+  worker
+extern void w_8 (void);
+#pragma acc routine (w_8) /* { dg-error "incompatible .seq. clause when applying .#pragma acc routine. to .\[void \]*w_8\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
+
+#pragma acc routine \
+  vector
+extern void v_8 (void);
+#pragma acc routine (v_8) /* { dg-error "incompatible .seq. clause when applying .#pragma acc routine. to .\[void \]*v_8\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
+
+extern void s_8 (void);
+#pragma acc routine (s_8)
+#pragma acc routine (s_8) \
+  vector /* { dg-error "incompatible .vector. clause when applying .#pragma acc routine. to .\[void \]*s_8\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
+#pragma acc routine (s_8) \
+  gang /* { dg-error "incompatible .gang. clause when applying .#pragma acc routine. to .\[void \]*s_8\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
+#pragma acc routine (s_8) \
+  worker /* { dg-error "incompatible .worker. clause when applying .#pragma acc routine. to .\[void \]*s_8\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-2.c b/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-2.c
new file mode 100644
index 0000000..1217397
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-2.c
@@ -0,0 +1,71 @@ 
+/* Test various aspects of clauses specifying compatible levels of
+   parallelism with the OpenACC routine directive.  The Fortran counterpart is
+   ../../gfortran.dg/goacc/routine-level-of-parallelism-2.f.  */
+
+#pragma acc routine gang
+void g_1 (void)
+{
+}
+#pragma acc routine (g_1) gang
+#pragma acc routine (g_1) gang
+
+
+extern void w_1 (void);
+#pragma acc routine (w_1) worker
+#pragma acc routine (w_1) worker
+#pragma acc routine (w_1) worker
+
+
+#pragma acc routine vector
+extern void v_1 (void);
+#pragma acc routine (v_1) vector
+#pragma acc routine (v_1) vector
+
+
+/* Also test the implicit seq clause.  */
+
+#pragma acc routine seq
+extern void s_1_1 (void);
+#pragma acc routine (s_1_1)
+#pragma acc routine (s_1_1) seq
+#pragma acc routine (s_1_1)
+#pragma acc routine (s_1_1) seq
+
+#pragma acc routine
+extern void s_1_2 (void);
+#pragma acc routine (s_1_2)
+#pragma acc routine (s_1_2) seq
+#pragma acc routine (s_1_2)
+#pragma acc routine (s_1_2) seq
+
+extern void s_2_1 (void);
+#pragma acc routine (s_2_1) seq
+#pragma acc routine (s_2_1)
+#pragma acc routine (s_2_1) seq
+#pragma acc routine (s_2_1)
+#pragma acc routine (s_2_1) seq
+
+extern void s_2_2 (void);
+#pragma acc routine (s_2_2)
+#pragma acc routine (s_2_2)
+#pragma acc routine (s_2_2) seq
+#pragma acc routine (s_2_2)
+#pragma acc routine (s_2_2) seq
+
+#pragma acc routine seq
+void s_3_1 (void)
+{
+}
+#pragma acc routine (s_3_1)
+#pragma acc routine (s_3_1) seq
+#pragma acc routine (s_3_1)
+#pragma acc routine (s_3_1) seq
+
+#pragma acc routine
+void s_3_2 (void)
+{
+}
+#pragma acc routine (s_3_2)
+#pragma acc routine (s_3_2) seq
+#pragma acc routine (s_3_2)
+#pragma acc routine (s_3_2) seq
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c b/gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c
new file mode 100644
index 0000000..17d6b03
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c
@@ -0,0 +1,48 @@ 
+/* Test the nohost clause for OpenACC routine directive.  Exercising different
+   variants for declaring routines.  */
+
+/* { dg-additional-options "-fdump-tree-oaccdevlow" } */
+
+#pragma acc routine nohost
+int THREE(void)
+{
+  return 3;
+}
+
+#pragma acc routine (THREE) nohost
+
+#pragma acc routine nohost
+extern int THREE(void);
+
+
+#pragma acc routine nohost
+extern void NOTHING(void);
+
+#pragma acc routine (NOTHING) nohost
+
+void NOTHING(void)
+{
+}
+
+#pragma acc routine nohost
+extern void NOTHING(void);
+
+#pragma acc routine (NOTHING) nohost
+
+
+extern float ADD(float, float);
+
+#pragma acc routine (ADD) nohost
+
+float ADD(float x, float y)
+{
+  return x + y;
+}
+
+#pragma acc routine nohost
+extern float ADD(float, float);
+
+#pragma acc routine (ADD) nohost
+
+
+/* { dg-final { scan-tree-dump-times "Discarding function" 3 "oaccdevlow" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-nohost-2.c b/gcc/testsuite/c-c++-common/goacc/routine-nohost-2.c
new file mode 100644
index 0000000..7402bfc
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/routine-nohost-2.c
@@ -0,0 +1,97 @@ 
+/* Test invalid usage of the nohost clause for OpenACC routine directive.
+   Exercising different variants for declaring routines.  */
+
+#pragma acc routine
+int THREE_1(void)
+{
+  return 3;
+}
+
+#pragma acc routine (THREE_1) \
+  nohost /* { dg-error "incompatible .nohost. clause when applying .#pragma acc routine. to .\[int \]*THREE_1\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
+
+#pragma acc routine \
+  nohost /* { dg-error "incompatible .nohost. clause when applying .#pragma acc routine. to .\[int \]*THREE_1\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
+extern int THREE_1(void);
+
+
+#pragma acc routine
+extern void NOTHING_1(void);
+
+#pragma acc routine (NOTHING_1) \
+  nohost /* { dg-error "incompatible .nohost. clause when applying .#pragma acc routine. to .\[void \]*NOTHING_1\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
+
+void NOTHING_1(void)
+{
+}
+
+#pragma acc routine \
+  nohost /* { dg-error "incompatible .nohost. clause when applying .#pragma acc routine. to .\[void \]*NOTHING_1\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
+extern void NOTHING_1(void);
+
+#pragma acc routine (NOTHING_1) \
+  nohost /* { dg-error "incompatible .nohost. clause when applying .#pragma acc routine. to .\[void \]*NOTHING_1\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
+
+
+extern float ADD_1(float, float);
+
+#pragma acc routine (ADD_1)
+
+float ADD_1(float x, float y)
+{
+  return x + y;
+}
+
+#pragma acc routine \
+  nohost /* { dg-error "incompatible .nohost. clause when applying .#pragma acc routine. to .\[float \]*ADD_1\[\\(float, \\)\]*., which has already been marked as an accelerator routine" } */
+extern float ADD_1(float, float);
+
+#pragma acc routine (ADD_1) \
+  nohost /* { dg-error "incompatible .nohost. clause when applying .#pragma acc routine. to .\[float \]*ADD_1\[\\(float, \\)\]*., which has already been marked as an accelerator routine" } */
+
+
+/* The same again, but with/without nohost reversed.  */
+
+#pragma acc routine \
+  nohost
+int THREE_2(void)
+{
+  return 3;
+}
+
+#pragma acc routine (THREE_2) /* { dg-error "missing .nohost. clause when applying .#pragma acc routine. to .\[int \]*THREE_2\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
+
+#pragma acc routine /* { dg-error "missing .nohost. clause when applying .#pragma acc routine. to .\[int \]*THREE_2\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
+extern int THREE_2(void);
+
+
+#pragma acc routine \
+  nohost
+extern void NOTHING_2(void);
+
+#pragma acc routine (NOTHING_2) /* { dg-error "missing .nohost. clause when applying .#pragma acc routine. to .\[void \]*NOTHING_2\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
+
+void NOTHING_2(void)
+{
+}
+
+#pragma acc routine /* { dg-error "missing .nohost. clause when applying .#pragma acc routine. to .\[void \]*NOTHING_2\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
+extern void NOTHING_2(void);
+
+#pragma acc routine (NOTHING_2) /* { dg-error "missing .nohost. clause when applying .#pragma acc routine. to .\[void \]*NOTHING_2\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
+
+
+extern float ADD_2(float, float);
+
+#pragma acc routine (ADD_2) \
+  nohost
+
+float ADD_2(float x, float y)
+{
+  return x + y;
+}
+
+#pragma acc routine /* { dg-error "missing .nohost. clause when applying .#pragma acc routine. to .\[float \]*ADD_2\[\\(float, \\)\]*., which has already been marked as an accelerator routine" } */
+extern float ADD_2(float, float);
+
+#pragma acc routine (ADD_2) /* { dg-error "missing .nohost. clause when applying .#pragma acc routine. to .\[float \]*ADD_2\[\\(float, \\)\]*., which has already been marked as an accelerator routine" } */
diff --git a/gcc/testsuite/gfortran.dg/goacc/fixed-1.f b/gcc/testsuite/gfortran.dg/goacc/fixed-1.f
index 6a454190..0c0fb98 100644
--- a/gcc/testsuite/gfortran.dg/goacc/fixed-1.f
+++ b/gcc/testsuite/gfortran.dg/goacc/fixed-1.f
@@ -1,3 +1,5 @@ 
+!$ACC ROUTINE(ABORT) SEQ
+
       INTEGER :: ARGC
       ARGC = COMMAND_ARGUMENT_COUNT ()
 
diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-5.f95 b/gcc/testsuite/gfortran.dg/goacc/loop-5.f95
index d059cf7..fe137d5 100644
--- a/gcc/testsuite/gfortran.dg/goacc/loop-5.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/loop-5.f95
@@ -93,9 +93,6 @@  program test
       DO j = 1,10
       ENDDO
     ENDDO
-    !$acc loop tile(-1) ! { dg-warning "must be positive" }
-    do i = 1,10
-    enddo
     !$acc loop vector tile(*)
     DO i = 1,10
     ENDDO
@@ -129,9 +126,6 @@  program test
       DO j = 1,10
       ENDDO
     ENDDO
-    !$acc loop tile(-1) ! { dg-warning "must be positive" }
-    do i = 1,10
-    enddo
     !$acc loop vector tile(*)
     DO i = 1,10
     ENDDO
@@ -242,9 +236,6 @@  program test
     DO j = 1,10
     ENDDO
   ENDDO
-  !$acc kernels loop tile(-1) ! { dg-warning "must be positive" }
-  do i = 1,10
-  enddo
   !$acc kernels loop vector tile(*)
   DO i = 1,10
   ENDDO
@@ -333,9 +324,6 @@  program test
     DO j = 1,10
     ENDDO
   ENDDO
-  !$acc parallel loop tile(-1) ! { dg-warning "must be positive" }
-  do i = 1,10
-  enddo
   !$acc parallel loop vector tile(*)
   DO i = 1,10
   ENDDO
diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-6.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-6.f90
index 10943cf..ee43935 100644
--- a/gcc/testsuite/gfortran.dg/goacc/routine-6.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-6.f90
@@ -5,7 +5,6 @@  contains
   subroutine subr5 (x) 
   implicit none
   !$acc routine (subr5)
-  !$acc routine (m1int) ! { dg-error "invalid function name" }
   integer, intent(inout) :: x
   if (x < 1) then
      x = 1
@@ -26,7 +25,6 @@  program main
   end interface
   integer, parameter :: n = 10
   integer :: a(n), i
-  !$acc routine (subr1) ! { dg-error "invalid function name" }
   external :: subr2
   !$acc routine (subr2)
 
@@ -56,7 +54,6 @@  subroutine subr1 (x)
 end subroutine subr1
 
 subroutine subr2 (x) 
-  !$acc routine (subr1) ! { dg-error "invalid function name" }
   integer, intent(inout) :: x
   if (x < 1) then
      x = 1
@@ -86,7 +83,6 @@  subroutine subr4 (x)
 end subroutine subr4
 
 subroutine subr10 (x)
-  !$acc routine (subr10) device ! { dg-error "Unclassifiable OpenACC directive" }
   integer, intent(inout) :: x
   if (x < 1) then
      x = 1
diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-7.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-7.f90
new file mode 100644
index 0000000..cb45079
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-7.f90
@@ -0,0 +1,70 @@ 
+! Test acc routines inside modules.
+
+! { dg-additional-options "-O0" }
+
+module routines
+  integer a
+contains
+  subroutine vector
+    implicit none
+    !$acc routine vector
+  end subroutine vector
+
+  subroutine worker
+    implicit none
+    !$acc routine worker
+  end subroutine worker
+
+  subroutine gang
+    implicit none
+    !$acc routine gang
+  end subroutine gang
+
+  subroutine seq
+    implicit none
+    !$acc routine seq
+  end subroutine seq
+end module routines
+
+program main
+  use routines
+  implicit none
+
+  integer :: i
+
+  !$acc parallel loop gang
+  do i = 1, 10
+     call gang ! { dg-error "routine call uses same OpenACC parallelism as containing loop" }
+     call worker
+     call vector
+     call seq
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop worker
+  do i = 1, 10
+     call gang ! { dg-error "routine call uses same OpenACC parallelism as containing loop" }
+     call worker ! { dg-error "routine call uses same OpenACC parallelism as containing loop" }
+     call vector
+     call seq
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop vector
+  do i = 1, 10
+     call gang ! { dg-error "routine call uses same OpenACC parallelism as containing loop" }
+     call worker ! { dg-error "routine call uses same OpenACC parallelism as containing loop" }
+     call vector ! { dg-error "routine call uses same OpenACC parallelism as containing loop" }
+     call seq
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop seq
+  do i = 1, 10
+     call gang
+     call worker
+     call vector
+     call seq
+  end do
+  !$acc end parallel loop
+end program main
diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-9.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-9.f90
new file mode 100644
index 0000000..590e594
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-9.f90
@@ -0,0 +1,96 @@ 
+! Check for late resolver errors caused by invalid ACC ROUTINE
+! directives.
+
+module m
+  integer m1int
+contains
+  subroutine subr5 (x)
+    implicit none
+    integer extfunc
+    !$acc routine (subr5)
+    !$acc routine (extfunc)
+    !$acc routine (m1int) ! { dg-error "invalid function name" }
+    integer, intent(inout) :: x
+    if (x < 1) then
+       x = 1
+    else
+       x = x * x - 1 + extfunc(2)
+    end if
+  end subroutine subr5
+end module m
+
+program main
+  implicit none
+  interface
+    function subr6 (x)
+    integer, intent (in) :: x
+    integer :: subr6
+    end function subr6
+  end interface
+  integer, parameter :: n = 10
+  integer :: a(n), i
+  !$acc routine (subr1) ! { dg-error "invalid function name" }
+  external :: subr2
+  !$acc routine (subr2)
+
+  external :: R1, R2
+  !$acc routine (R1)
+  !$acc routine (R2)
+
+  !$acc parallel
+  !$acc loop
+  do i = 1, n
+     call subr1 (i)
+     call subr2 (i)
+  end do
+  !$acc end parallel
+end program main
+
+subroutine subr1 (x)
+  !$acc routine
+  integer, intent(inout) :: x
+  if (x < 1) then
+     x = 1
+  else
+     x = x * x - 1
+  end if
+end subroutine subr1
+
+subroutine subr2 (x)
+  integer, intent(inout) :: x
+  if (x < 1) then
+     x = 1
+  else
+     x = x * x - 1
+  end if
+end subroutine subr2
+
+subroutine subr3 (x)
+  !$acc routine (subr3)
+  integer, intent(inout) :: x
+  if (x < 1) then
+     x = 1
+  else
+     call subr4 (x)
+  end if
+end subroutine subr3
+
+subroutine subr4 (x)
+  !$acc routine (subr4)
+  integer, intent(inout) :: x
+  if (x < 1) then
+     x = 1
+  else
+     x = x * x - 1
+  end if
+end subroutine subr4
+
+subroutine subr10 (x)
+  !$acc routine (subr10) device ! { dg-error "Unclassifiable OpenACC directive" }
+  integer, intent(inout) :: x
+  if (x < 1) then
+     x = 1
+  else
+     x = x * x - 1
+  end if
+end subroutine subr10
diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-nested-parallelism.f b/gcc/testsuite/gfortran.dg/goacc/routine-nested-parallelism.f
new file mode 100644
index 0000000..d1304c6
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-nested-parallelism.f
@@ -0,0 +1,340 @@ 
+! Validate calls to ACC ROUTINES.  Ensure that the loop containing the
+! call has sufficient parallelism to for the routine.
+
+      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
+      do i = 1, n
+         !$acc loop
+         do j = 1, n
+            call workerr (a, n)
+         end do
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop
+      do i = 1, n
+!$acc loop gang
+         do j = 1, n
+            call gangr (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" }
+         end do
+      end do
+!$acc end parallel loop
+
+!
+! Test calls to seq routines
+!
+
+!$acc parallel loop
+      do i = 1, n
+         call seqr (a, n)
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop gang
+      do i = 1, n
+         call seqr (a, n)
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop worker
+      do i = 1, n
+         call seqr (a, n)
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop vector
+      do i = 1, n
+         call seqr (a, n)
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop seq
+      do i = 1, n
+         call seqr (a, n)
+      end do
+!$acc end parallel loop
+
+!
+! Test calls to gang routines
+!
+
+!$acc parallel loop
+      do i = 1, n
+         call gangr (a, n)
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop gang
+      do i = 1, n
+         call gangr (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop worker
+      do i = 1, n
+         call gangr (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop vector
+      do i = 1, n
+         call gangr (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop seq
+      do i = 1, n
+         call gangr (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" }
+      end do
+!$acc end parallel loop
+
+!
+! Test calls to worker routines
+!
+
+!$acc parallel loop
+      do i = 1, n
+         call workerr (a, n)
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop gang
+      do i = 1, n
+         call workerr (a, n)
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop worker
+      do i = 1, n
+         call workerr (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop vector
+      do i = 1, n
+         call workerr (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop seq
+      do i = 1, n
+         call workerr (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" }
+      end do
+!$acc end parallel loop
+
+!
+! Test calls to vector routines
+!
+
+!$acc parallel loop
+      do i = 1, n
+         call vectorr (a, n)
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop gang
+      do i = 1, n
+         call vectorr (a, n)
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop worker
+      do i = 1, n
+         call vectorr (a, n)
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop vector
+      do i = 1, n
+         call vectorr (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop seq
+      do i = 1, n
+         call vectorr (a, n) ! { dg-error "Insufficient ..ACC 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 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
+      do i = 1, n
+!$acc loop
+         do j = 1, n
+            a(1) = workerf (a, n)
+         end do
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop
+      do i = 1, n
+!$acc loop gang
+         do j = 1, n
+            a(1) = gangf (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" }
+         end do
+      end do
+!$acc end parallel loop
+
+!
+! Test calls to seq routines
+!
+
+!$acc parallel loop
+      do i = 1, n
+         a(1) = seqf (a, n)
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop gang
+      do i = 1, n
+         a(1) = seqf (a, n)
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop worker
+      do i = 1, n
+         a(1) = seqf (a, n)
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop vector
+      do i = 1, n
+         a(1) = seqf (a, n)
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop seq
+      do i = 1, n
+         a(1) = seqf (a, n)
+      end do
+!$acc end parallel loop
+
+!
+! Test calls to gang routines
+!
+
+!$acc parallel loop
+      do i = 1, n
+         a(1) = gangf (a, n)
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop gang
+      do i = 1, n
+         a(1) = gangf (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop worker
+      do i = 1, n
+         a(1) = gangf (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop vector
+      do i = 1, n
+         a(1) = gangf (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop seq
+      do i = 1, n
+         a(1) = gangf (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" }
+      end do
+!$acc end parallel loop
+
+!
+! Test calls to worker routines
+!
+
+!$acc parallel loop
+      do i = 1, n
+         a(1) = workerf (a, n)
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop gang
+      do i = 1, n
+         a(1) = workerf (a, n)
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop worker
+      do i = 1, n
+         a(1) = workerf (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop vector
+      do i = 1, n
+         a(1) = workerf (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop seq
+      do i = 1, n
+         a(1) = workerf (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" }
+      end do
+!$acc end parallel loop
+
+!
+! Test calls to vector routines
+!
+
+!$acc parallel loop
+      do i = 1, n
+         a(1) = vectorf (a, n)
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop gang
+      do i = 1, n
+         a(1) = vectorf (a, n)
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop worker
+      do i = 1, n
+         a(1) = vectorf (a, n)
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop vector
+      do i = 1, n
+         a(1) = vectorf (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" }
+      end do
+!$acc end parallel loop
+
+!$acc parallel loop seq
+      do i = 1, n
+         a(1) = vectorf (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" }
+      end do
+!$acc end parallel loop
+      end subroutine func
diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-nested-parallelism.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-nested-parallelism.f90
new file mode 100644
index 0000000..94e0464
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-nested-parallelism.f90
@@ -0,0 +1,340 @@ 
+! Validate calls to ACC ROUTINES.  Ensure that the loop containing the
+! call has sufficient parallelism to for the routine.
+
+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
+  do i = 1, n
+     !$acc loop
+     do j = 1, n
+        call workerr (a, n)
+     end do
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop
+  do i = 1, n
+     !$acc loop gang
+     do j = 1, n
+        call gangr (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" }
+     end do
+  end do
+  !$acc end parallel loop
+
+  !
+  ! Test calls to seq routines
+  !
+
+  !$acc parallel loop
+  do i = 1, n
+     call seqr (a, n)
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop gang
+  do i = 1, n
+     call seqr (a, n)
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop worker
+  do i = 1, n
+     call seqr (a, n)
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop vector
+  do i = 1, n
+     call seqr (a, n)
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop seq
+  do i = 1, n
+     call seqr (a, n)
+  end do
+  !$acc end parallel loop
+
+  !
+  ! Test calls to gang routines
+  !
+
+  !$acc parallel loop
+  do i = 1, n
+     call gangr (a, n)
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop gang
+  do i = 1, n
+     call gangr (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" }
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop worker
+  do i = 1, n
+     call gangr (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" }
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop vector
+  do i = 1, n
+     call gangr (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" }
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop seq
+  do i = 1, n
+     call gangr (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" }
+  end do
+  !$acc end parallel loop
+
+  !
+  ! Test calls to worker routines
+  !
+
+  !$acc parallel loop
+  do i = 1, n
+     call workerr (a, n)
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop gang
+  do i = 1, n
+     call workerr (a, n)
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop worker
+  do i = 1, n
+     call workerr (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" }
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop vector
+  do i = 1, n
+     call workerr (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" }
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop seq
+  do i = 1, n
+     call workerr (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" }
+  end do
+  !$acc end parallel loop
+
+  !
+  ! Test calls to vector routines
+  !
+
+  !$acc parallel loop
+  do i = 1, n
+     call vectorr (a, n)
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop gang
+  do i = 1, n
+     call vectorr (a, n)
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop worker
+  do i = 1, n
+     call vectorr (a, n)
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop vector
+  do i = 1, n
+     call vectorr (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" }
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop seq
+  do i = 1, n
+     call vectorr (a, n) ! { dg-error "Insufficient ..ACC 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 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
+  do i = 1, n
+     !$acc loop
+     do j = 1, n
+        a(1) = workerf (a, n)
+     end do
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop
+  do i = 1, n
+     !$acc loop gang
+     do j = 1, n
+        a(1) = gangf (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" }
+     end do
+  end do
+  !$acc end parallel loop
+
+  !
+  ! Test calls to seq routines
+  !
+
+  !$acc parallel loop
+  do i = 1, n
+     a(1) = seqf (a, n)
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop gang
+  do i = 1, n
+     a(1) = seqf (a, n)
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop worker
+  do i = 1, n
+     a(1) = seqf (a, n)
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop vector
+  do i = 1, n
+     a(1) = seqf (a, n)
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop seq
+  do i = 1, n
+     a(1) = seqf (a, n)
+  end do
+  !$acc end parallel loop
+
+  !
+  ! Test calls to gang routines
+  !
+
+  !$acc parallel loop
+  do i = 1, n
+     a(1) = gangf (a, n)
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop gang
+  do i = 1, n
+     a(1) = gangf (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" }
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop worker
+  do i = 1, n
+     a(1) = gangf (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" }
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop vector
+  do i = 1, n
+     a(1) = gangf (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" }
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop seq
+  do i = 1, n
+     a(1) = gangf (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" }
+  end do
+  !$acc end parallel loop
+
+  !
+  ! Test calls to worker routines
+  !
+
+  !$acc parallel loop
+  do i = 1, n
+     a(1) = workerf (a, n)
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop gang
+  do i = 1, n
+     a(1) = workerf (a, n)
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop worker
+  do i = 1, n
+     a(1) = workerf (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" }
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop vector
+  do i = 1, n
+     a(1) = workerf (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" }
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop seq
+  do i = 1, n
+     a(1) = workerf (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" }
+  end do
+  !$acc end parallel loop
+
+  !
+  ! Test calls to vector routines
+  !
+
+  !$acc parallel loop
+  do i = 1, n
+     a(1) = vectorf (a, n)
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop gang
+  do i = 1, n
+     a(1) = vectorf (a, n)
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop worker
+  do i = 1, n
+     a(1) = vectorf (a, n)
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop vector
+  do i = 1, n
+     a(1) = vectorf (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" }
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop seq
+  do i = 1, n
+     a(1) = vectorf (a, n) ! { dg-error "Insufficient ..ACC LOOP parallelism" }
+  end do
+  !$acc end parallel loop
+end subroutine func
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-3.c
new file mode 100644
index 0000000..2cdd6bf
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-3.c
@@ -0,0 +1,33 @@ 
+/* At -O0, we do get the expected "undefined reference to `foo'" link-time
+   error message (but the check needs to be done differently; compare to
+   routine-nohost-1.c), but for -O2 we don't; presumably because the function
+   gets inlined.
+   { dg-xfail-if "TODO" { *-*-* } { "-O0" } { "" } } */
+
+#include <stdlib.h>
+
+#pragma acc routine nohost
+int
+foo (int n)
+{
+  if (n == 0 || n == 1)
+    return 1;
+
+  return n * n;
+}
+
+int
+main()
+{
+  int a, n = 10;
+
+#pragma acc parallel copy (a, n)
+  {
+    a = foo (n);
+  }
+
+  if (a != n * n)
+    abort ();
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-5.c
new file mode 100644
index 0000000..4e34f78
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-5.c
@@ -0,0 +1,64 @@ 
+
+/* { dg-do run } */
+/* { dg-warning "TODO" "implicit" { xfail *-*-* } 17 } */
+/* { dg-warning "TODO" "implicit" { xfail *-*-* } 27 } */
+/* { dg-xfail-if "unresolved symbol" { *-*-* } } */
+
+#include <stdlib.h>
+
+#pragma acc routine bind (foo)
+int
+subr1 (int n)
+{
+  if (n == 0 || n == 1)
+    return 1;
+
+  return n * foo (n - 1);
+}
+
+#pragma acc routine bind ("bar")
+int
+subr2 (int n)
+{
+  if (n == 0 || n == 1)
+    return 1;
+
+  return n * bar (n - 1);
+}
+
+int
+main()
+{
+  int *a, i, n = 10;
+
+  a = (int *)malloc (sizeof (int) * n);
+
+#pragma acc parallel copy (a[0:n]) vector_length (5)
+  {
+#pragma acc loop
+    for (i = 0; i < n; i++)
+      a[i] = foo (i);
+  }
+
+  for (i = 0; i < n; i++)
+    if (a[i] != subr1 (i))
+      abort ();
+
+  for (i = 0; i < n; i++)
+    a[i] = 0;
+
+#pragma acc parallel copy (a[0:n]) vector_length (5)
+  {
+#pragma acc loop
+    for (i = 0; i < n; i++)
+      a[i] = bar (i);
+  }
+
+  for (i = 0; i < n; i++)
+    if (a[i] != subr2 (i))
+      abort ();
+
+  free (a);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-bind-nohost-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-bind-nohost-1.c
new file mode 100644
index 0000000..b991bb1
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-bind-nohost-1.c
@@ -0,0 +1,105 @@ 
+/* Test the bind and nohost clauses for OpenACC routine directive.  */
+
+/* TODO.  Function inlining and the OpenACC bind clause do not yet get on well
+   with one another.
+   { dg-additional-options "-fno-inline" } */
+
+/* TODO.  C works, but for C++ we get: "lto1: internal compiler error: in
+   ipa_propagate_frequency".
+   { dg-xfail-if "TODO" { *-*-* } } */
+
+#include <openacc.h>
+
+/* "MINUS_TWO" is the device variant for function "TWO".  Similar for "THREE",
+   and "FOUR".  Exercising different variants for declaring routines.  */
+
+#pragma acc routine nohost
+extern int MINUS_TWO(void);
+
+int MINUS_TWO(void)
+{
+  if (!acc_on_device(acc_device_not_host))
+    __builtin_abort();
+  return -2;
+}
+
+extern int TWO(void);
+#pragma acc routine (TWO) bind(MINUS_TWO)
+
+int TWO(void)
+{
+  if (acc_on_device(acc_device_not_host))
+    __builtin_abort();
+  return 2;
+}
+
+
+#pragma acc routine nohost
+int MINUS_THREE(void)
+{
+  if (!acc_on_device(acc_device_not_host))
+    __builtin_abort();
+  return -3;
+}
+
+#pragma acc routine bind(MINUS_THREE)
+extern int THREE(void);
+
+int THREE(void)
+{
+  if (acc_on_device(acc_device_not_host))
+    __builtin_abort();
+  return 3;
+}
+
+
+/* Due to using a string in the bind clause, we don't need "MINUS_FOUR" in
+   scope here.  */
+#pragma acc routine bind("MINUS_FOUR")
+int FOUR(void)
+{
+  if (acc_on_device(acc_device_not_host))
+    __builtin_abort();
+  return 4;
+}
+
+extern int MINUS_FOUR(void);
+#pragma acc routine (MINUS_FOUR) nohost
+
+int MINUS_FOUR(void)
+{
+  if (!acc_on_device(acc_device_not_host))
+    __builtin_abort();
+  return -4;
+}
+
+
+int main()
+{
+  int x2, x3, x4;
+
+#pragma acc parallel copyout(x2, x3, x4) if(0)
+  {
+    x2 = TWO();
+    x3 = THREE();
+    x4 = FOUR();
+  }
+  if (x2 != 2 || x3 != 3 || x4 != 4)
+    __builtin_abort();
+
+#pragma acc parallel copyout(x2, x3, x4)
+  {
+    x2 = TWO();
+    x3 = THREE();
+    x4 = FOUR();
+  }
+#ifdef ACC_DEVICE_TYPE_host
+  if (x2 != 2 || x3 != 3 || x4 != 4)
+    __builtin_abort();
+#else
+  if (x2 != -2 || x3 != -3 || x4 != -4)
+    __builtin_abort();
+#endif
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-nohost-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-nohost-1.c
new file mode 100644
index 0000000..365af93
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-nohost-1.c
@@ -0,0 +1,18 @@ 
+/* { dg-do link } */
+
+extern int three (void);
+
+#pragma acc routine (three) nohost
+__attribute__((noinline))
+int three(void)
+{
+  return 3;
+}
+
+int main(void)
+{
+  return (three() == 3) ? 0 : 1;
+}
+
+/* Expecting link to fail; "undefined reference to `three'" (or similar).
+   { dg-excess-errors "" } */
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/abort-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/abort-1.f90
index b38303d..48ebc38 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/abort-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/abort-1.f90
@@ -1,5 +1,6 @@ 
 program main
   implicit none
+  !$acc routine(abort) seq
 
   print *, "CheCKpOInT"
   !$acc parallel
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f b/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f
index a19045b..cbd1dd9 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f
+++ b/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f
@@ -6,6 +6,7 @@ 
 
       USE OPENACC
       IMPLICIT NONE
+!$ACC ROUTINE(ABORT) SEQ
 
 !Host.
 
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/routine-6.f90 b/libgomp/testsuite/libgomp.oacc-fortran/routine-6.f90
new file mode 100644
index 0000000..1bae09c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/routine-6.f90
@@ -0,0 +1,28 @@ 
+! { dg-do run }
+! { dg-xfail-if "TODO" { *-*-* } }
+
+program main
+  integer :: a, n
+  
+  n = 10
+
+  !$acc parallel copy (a, n)
+     a = foo (n)
+  !$acc end parallel 
+
+  if (a .ne. n * n) call abort
+
+contains
+
+function foo (n) result (rc)
+  !$acc routine nohost
+
+  integer, intent (in) :: n
+  integer :: rc
+
+  rc = n * n
+
+end function
+
+end program main
+
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/routine-7.f90 b/libgomp/testsuite/libgomp.oacc-fortran/routine-7.f90
index 200188e..ef2ff04 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/routine-7.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/routine-7.f90
@@ -100,7 +100,7 @@  subroutine gang (a)
   integer, intent (inout) :: a(N)
   integer :: i
 
-  !$acc loop gang
+  !$acc loop gang worker vector
   do i = 1, N
     a(i) = a(i) - i 
   end do
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/routine-8.f90 b/libgomp/testsuite/libgomp.oacc-fortran/routine-8.f90
new file mode 100644
index 0000000..5c58b43
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/routine-8.f90
@@ -0,0 +1,62 @@ 
+
+! { dg-do run } 
+! { dg-error "Invalid" "TODO" { xfail *-*-* } 51 }
+
+program main
+  integer, parameter :: n = 10
+  integer :: a(n)
+  integer :: i
+
+  !$acc parallel copy (a) vector_length (5)
+  !$acc loop
+    do i = 1, n
+      a(i) = foo (i);
+    end do
+  !$acc end parallel
+
+  do i = 1, n
+    if (a(i) .ne. subr1 (i)) call abort
+  end do
+
+  do i = 1, n
+    a(i) = 0
+  end do
+
+  !$acc parallel copy (a) vector_length (5)
+  !$acc loop
+    do i = 1, n
+      a(i) = bar (i);
+    end do
+  !$acc end parallel
+
+  do i = 1, n
+    if (a(i) .ne. subr2 (i)) call abort
+  end do
+
+contains
+
+function subr1 (n) result (rc)
+  !$acc routine bind (foo)
+  integer :: n, rc
+
+  if ((n .eq. 0) .or. (n .eq. 1)) then
+    rc = 1
+  else
+    rc = n * foo (n - 1);
+  end if
+
+end function
+
+function subr2 (n) result (rc)
+  !$acc routine bind ("bar")
+  integer :: n, rc
+
+  if ((n .eq. 0) .or. (n .eq. 1)) then
+    rc = 1
+  else
+    rc = n * bar (n - 1);
+  end if
+
+end function
+
+end program main
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/vector-routine.f90 b/libgomp/testsuite/libgomp.oacc-fortran/vector-routine.f90
new file mode 100644
index 0000000..1edcee4
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/vector-routine.f90
@@ -0,0 +1,41 @@ 
+! { dg-do run }
+
+module param
+  integer, parameter :: N = 32
+end module param
+
+program main
+  use param
+  integer :: i
+  integer :: a(N)
+
+  do i = 1, N
+    a(i) = i
+  end do
+
+  !$acc parallel copy (a)
+  !$acc loop worker
+    do i = 1, N
+      call vector (a)
+    end do
+  !$acc end parallel
+
+  do i = 1, N
+    if (a(i) .ne. 0) call abort
+  end do
+
+contains
+
+  subroutine vector (a)
+  !$acc routine vector
+  integer, intent (inout) :: a(N)
+  integer :: i
+
+  !$acc loop vector
+  do i = 1, N
+    a(i) = a(i) - a(i) 
+  end do
+
+end subroutine vector
+
+end program main