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

