This is the mail archive of the gcc-patches@gcc.gnu.org mailing list for the GCC project.


Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]
Other format: [Raw text]

[gomp4] Make front ends emit warnings when OpenACC routines directives lack parallelism


The language regarding OpenACC routines have changed in OpenACC 2.5 such
that the end user must explicitly specify one of gang, worker, vector or
seq partitioning. I guess some other compiler need those directives to
generate specialized versions of those functions accordingly. However,
GCC currently falls back to 'seq' partitioning, and that seems to be a
reasonable compromise in most cases.

This patch teaches the FE's how to emit a warning whenever a routine
directive doesn't include a partitioning clause. The rationale for
making this a warning vs an error is that it enables us to have some
backwards compatibility with OpenACC 2.0a.

Unfortunately, the fortran FE doesn't make use of
verify_oacc_routine_clauses because it parses/matches routines much
earlier than c/c++, so that part of the patch is slightly more involved.
I've also had to update a lot of test cases to make them conform to the
new routine behavior.

This patch has been committed to gomp-4_0-branch.

Cesar
2017-05-01  Cesar Philippidis  <cesar@codesourcery.com>

	gcc/fortran/
	* gfortran.h (enum oacc_function): Add OACC_FUNCTION_AUTO.
	* openmp.c (gfc_oacc_routine_dims): Return OACC_FUNCTION_AUTO when no
	parallelism was detected.
	(gfc_match_oacc_routine): Emit a warning when the user doesn't
	supply a gang, worker, vector or seq clause to an OpenACC routine
	construct.

	gcc/
	* omp-low.c (verify_oacc_routine_clauses): Emit a warning when the
	user doesn't supply a gang, worker, vector or seq clause to an
	OpenACC routine construct.

	gcc/testsuite/
	* c-c++-common/goacc-gomp/nesting-fail-1.c: Update usage and expected
	output of OpenACC routines.
	* c-c++-common/goacc/Wparentheses-1.c: Likewise.
	* c-c++-common/goacc/nesting-fail-1.c: Likewise.
	* c-c++-common/goacc/routine-1.c: Likewise.
	* 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: Likewise.
	* c-c++-common/goacc/routine-level-of-parallelism-2.c: Likewise.
	* c-c++-common/goacc/routine-nohost-1.c: Likewise.
	* c-c++-common/goacc/routine-nohost-2.c: Likewise.
	* g++.dg/goacc/routine-1.C: Likewise.
	* g++.dg/goacc/routine-2.C: Likewise.
	* g++.dg/goacc/template.C: Likewise.
	* gfortran.dg/goacc/dtype-1.f95: Likewise.
	* gfortran.dg/goacc/pr71704.f90: Likewise.
	* gfortran.dg/goacc/pr72741-2.f: Likewise.
	* gfortran.dg/goacc/routine-6.f90: Likewise.
	* gfortran.dg/goacc/routine-8.f90: Likewise.
	* gfortran.dg/goacc/routine-9.f90: Likewise.
	* gfortran.dg/goacc/routine-level-of-parallelism-1.f90: Likewise.
	* gfortran.dg/goacc/routine-without-clauses.f90: New test.

	libgomp/
	* testsuite/libgomp.oacc-c++/pr71959-a.C: Adjust test case conform
	to OpenACC 2.5 routines. 
	* testsuite/libgomp.oacc-c-c++-common/declare-2.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/declare-3.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/declare-4.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/host_data-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/loop-default.h: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/mode-transitions.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/parallel-loop-2.h: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/routine-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/routine-3.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/routine-bind-nohost-1.c:
	Likewise.
	* testsuite/libgomp.oacc-fortran/host_data-2.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/host_data-3.f: Likewise.
	* testsuite/libgomp.oacc-fortran/host_data-4.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/routine-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/routine-2.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/routine-3.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/routine-4.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/routine-5.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/routine-9.f90: Likewise.

diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index 7913ac7..3c762a8 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -316,7 +316,8 @@ enum oacc_function
   OACC_FUNCTION_GANG,
   OACC_FUNCTION_WORKER,
   OACC_FUNCTION_VECTOR,
-  OACC_FUNCTION_SEQ
+  OACC_FUNCTION_SEQ,
+  OACC_FUNCTION_AUTO
 };
 
 /* Strings for all symbol attributes.  We use these for dumping the
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index 72c6669..88ccff2 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -2379,7 +2379,7 @@ static oacc_function
 gfc_oacc_routine_dims (gfc_omp_clauses *clauses)
 {
   int level = -1;
-  oacc_function ret = OACC_FUNCTION_SEQ;
+  oacc_function ret = OACC_FUNCTION_AUTO;
 
   if (clauses)
     {
@@ -2401,7 +2401,10 @@ gfc_oacc_routine_dims (gfc_omp_clauses *clauses)
 	  ret = OACC_FUNCTION_VECTOR;
 	}
       if (clauses->seq)
-	level = GOMP_DIM_MAX, mask |= GOMP_DIM_MASK (level);
+	{
+	  level = GOMP_DIM_MAX, mask |= GOMP_DIM_MASK (level);
+	  ret = OACC_FUNCTION_SEQ;
+	}
 
       if (mask != (mask & -mask))
 	ret = OACC_FUNCTION_NONE;
@@ -2505,6 +2508,12 @@ gfc_match_oacc_routine (void)
 	 know of any potential duplicate routine directives.  */
       seen_error = true;
     }
+  else if (dims == OACC_FUNCTION_AUTO)
+    {
+      gfc_warning (0, "Expected one of %<gang%>, %<worker%>, %<vector%> or "
+		   "%<seq%> clauses in !$ACC ROUTINE at %L", &old_loc);
+      dims = OACC_FUNCTION_SEQ;
+    }
 
   if (isym != NULL)
     {
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index cc209ba..88a1bb9 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -13239,8 +13239,10 @@ verify_oacc_routine_clauses (tree fndecl, tree *clauses, location_t loc,
       }
   if (c_level == NULL_TREE)
     {
-      /* OpenACC 2.5 makes this an error; for the current OpenACC 2.0a
-	 implementation add an implicit "seq" clause.  */
+      /* OpenACC 2.5 expects the user to supply one parallelism clause.  */
+      warning_at (loc, 0, "expecting one of %<gang%>, %<worker%>, %<vector%> "
+		  "or %<seq%> clauses");
+      inform (loc, "assigning %<seq%> parallelism to this routine");
       c_level = build_omp_clause (loc, OMP_CLAUSE_SEQ);
       OMP_CLAUSE_CHAIN (c_level) = *clauses;
       *clauses = c_level;
diff --git a/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c b/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
index 1a33242..57eaa02 100644
--- a/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
+++ b/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
@@ -362,7 +362,7 @@ f_acc_data (void)
   }
 }
 
-#pragma acc routine
+#pragma acc routine seq
 void
 f_acc_loop (void)
 {
@@ -436,7 +436,7 @@ f_acc_loop (void)
     }
 }
 
-#pragma acc routine
+#pragma acc routine seq
 void
 f_acc_routine (void)
 {
diff --git a/gcc/testsuite/c-c++-common/goacc/Wparentheses-1.c b/gcc/testsuite/c-c++-common/goacc/Wparentheses-1.c
index 08265b6..9d35ec5 100644
--- a/gcc/testsuite/c-c++-common/goacc/Wparentheses-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/Wparentheses-1.c
@@ -4,9 +4,9 @@
 int a, b, c;
 void bar (void);
 void baz (void);
-#pragma acc routine
+#pragma acc routine seq
 void bar2 (void);
-#pragma acc routine
+#pragma acc routine seq
 void baz2 (void);
 
 void
diff --git a/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c b/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c
index 93a9111..1462095 100644
--- a/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c
@@ -57,7 +57,7 @@ f_acc_data (void)
   }
 }
 
-#pragma acc routine
+#pragma acc routine seq
 void
 f_acc_routine (void)
 {
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-1.c b/gcc/testsuite/c-c++-common/goacc/routine-1.c
index 0b56661..a4ecfd3 100644
--- a/gcc/testsuite/c-c++-common/goacc/routine-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/routine-1.c
@@ -21,15 +21,15 @@ void seq (void)
 }
 
 
-#pragma acc routine
+#pragma acc routine /* { dg-warning "expecting one of" } */
 void bind_f_1 (void)
 {
 }
 
-#pragma acc routine
+#pragma acc routine /* { dg-warning "expecting one of" } */
 extern void bind_f_1 (void);
 
-#pragma acc routine (bind_f_1)
+#pragma acc routine (bind_f_1)/* { dg-warning "expecting one of" } */
 
 
 #pragma acc routine \
@@ -106,3 +106,16 @@ int main ()
 
   return 0;
 }
+
+/* { dg-warning "expecting one of" "" { target *-*-* } 35 } */
+/* { dg-warning "expecting one of" "" { target *-*-* } 41 } */
+/* { dg-warning "expecting one of" "" { target *-*-* } 45 } */
+/* { dg-warning "expecting one of" "" { target *-*-* } 50 } */
+/* { dg-warning "expecting one of" "" { target *-*-* } 56 } */
+/* { dg-warning "expecting one of" "" { target *-*-* } 60 } */
+/* { dg-warning "expecting one of" "" { target *-*-* } 64 } */
+/* { dg-warning "expecting one of" "" { target *-*-* } 70 } */
+/* { dg-warning "expecting one of" "" { target *-*-* } 74 } */
+/* { dg-warning "expecting one of" "" { target *-*-* } 78 } */
+/* { dg-warning "expecting one of" "" { target *-*-* } 84 } */
+/* { dg-warning "expecting one of" "" { target *-*-* } 88 } */
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-2.c b/gcc/testsuite/c-c++-common/goacc/routine-2.c
index debf6d7..65c7963 100644
--- a/gcc/testsuite/c-c++-common/goacc/routine-2.c
+++ b/gcc/testsuite/c-c++-common/goacc/routine-2.c
@@ -1,171 +1,171 @@
 /* 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" } */
+#pragma acc routine seq 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" } */
+#pragma acc routine seq 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" }*/
+#pragma acc routine seq 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" } */
+#pragma acc routine seq 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" } */
+#pragma acc routine seq 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" } */
+#pragma acc routine seq 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" } */
+#pragma acc routine seq 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 (a_bind_f_1) seq
 
 
 #pragma acc routine \
-  bind (a_bind_f_1)
+  bind (a_bind_f_1) seq
 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" } */
+#pragma acc routine seq /* { 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" } */
+#pragma acc routine (a_bind_f_1_1) seq /* { 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")
+  bind ("a_bind_f_2") seq
 void a_bind_f_2 (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" } */
+#pragma acc routine seq /* { 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 (a_bind_f_2) seq /* { 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")
+  bind ("a_bind_f_2") seq
 void a_bind_f_2_1 (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" } */
+#pragma acc routine seq /* { 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" } */
+#pragma acc routine (a_bind_f_2_1) seq /* { 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
+#pragma acc routine seq
 extern void b_bind_f_1 (void);
 
 
-#pragma acc routine
+#pragma acc routine seq
 void b_bind_f_1_1 (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" } */
+  bind (b_bind_f_1) seq /* { 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" } */
+  bind (b_bind_f_1) seq /* { 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
+#pragma acc routine seq
 void b_bind_f_2 (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\[\\(\\)\]*., which has already been marked as an accelerator routine" } */
+  bind ("b_bind_f_2") seq /* { 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" } */
+  bind ("b_bind_f_2") seq /* { 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
+#pragma acc routine seq
 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" } */
+  bind ("b_bind_f_2") seq /* { 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" } */
+  bind ("b_bind_f_2") seq /* { 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
+#pragma acc routine seq
 void c_bind_f_1a (void)
 {
 }
 
-#pragma acc routine
+#pragma acc routine seq
 extern void c_bind_f_1b (void);
 
 
 #pragma acc routine \
-  bind (c_bind_f_1a)
+  bind (c_bind_f_1a) seq
 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" } */
+  bind (c_bind_f_1b) seq /* { 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" } */
+  bind (c_bind_f_1b) seq /* { 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")
+  bind ("c_bind_f_2") seq
 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" } */
+  bind ("C_BIND_F_2") seq /* { 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" } */
+  bind ("C_BIND_F_2") seq /* { 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")
+  bind ("c_bind_f_2") seq
 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" } */
+  bind ("C_BIND_F_2") seq /* { 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" } */
+  bind ("C_BIND_F_2") seq /* { 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 f10651d..b7e8402 100644
--- a/gcc/testsuite/c-c++-common/goacc/routine-5.c
+++ b/gcc/testsuite/c-c++-common/goacc/routine-5.c
@@ -4,11 +4,11 @@
 
 struct PC
 {
-#pragma acc routine /* { dg-error ".#pragma acc routine. must be at file scope" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. must be at file scope" } */
 };
 
 void PC1( /* { dg-bogus "variable or field .PC1. declared void" "TODO" { xfail c++ } } */
-#pragma acc routine
+#pragma acc routine seq
 	 /* { dg-error ".#pragma acc routine. must be at file scope" "" { target c } 11 }
 	    { dg-error ".#pragma. is not allowed here" "" { target c++ } 11 } */
 ) /* { dg-bogus "expected declaration specifiers or .\\.\\.\\.. before .\\). token" "TODO" { xfail c } } */
@@ -18,26 +18,26 @@ void PC1( /* { dg-bogus "variable or field .PC1. declared void" "TODO" { xfail c
 void PC2()
 {
   if (0)
-#pragma acc routine /* { dg-error ".#pragma acc routine. must be at file scope" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. must be at file scope" } */
     ;
 }
 
 void PC3()
 {
-#pragma acc routine /* { dg-error ".#pragma acc routine. must be at file scope" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. must be at file scope" } */
 }
 
 
 /* "( name )" syntax.  */
 
 #pragma acc routine ( /* { dg-error "expected (function name|unqualified-id) before end of line" } */
-#pragma acc routine () /* { dg-error "expected (function name|unqualified-id) before .\\). token" } */
-#pragma acc routine (+) /* { dg-error "expected (function name|unqualified-id) before .\\+. token" } */
-#pragma acc routine (?) /* { dg-error "expected (function name|unqualified-id) before .\\?. token" } */
-#pragma acc routine (:) /* { dg-error "expected (function name|unqualified-id) before .:. token" } */
-#pragma acc routine (4) /* { dg-error "expected (function name|unqualified-id) before numeric constant" } */
+#pragma acc routine () seq /* { dg-error "expected (function name|unqualified-id) before .\\). token" } */
+#pragma acc routine (+) seq /* { dg-error "expected (function name|unqualified-id) before .\\+. token" } */
+#pragma acc routine (?) seq /* { dg-error "expected (function name|unqualified-id) before .\\?. token" } */
+#pragma acc routine (:) seq /* { dg-error "expected (function name|unqualified-id) before .:. token" } */
+#pragma acc routine (4) seq /* { dg-error "expected (function name|unqualified-id) before numeric constant" } */
 #pragma acc routine ('4') /* { dg-error "expected (function name|unqualified-id) before .4." } */
-#pragma acc routine ("4") /* { dg-error "expected (function name|unqualified-id) before string constant" } */
+#pragma acc routine ("4") seq /* { dg-error "expected (function name|unqualified-id) before string constant" } */
 extern void R1(void);
 extern void R2(void);
 #pragma acc routine (R1, R2, R3) worker /* { dg-error "expected .\\). before .,. token" } */
@@ -49,84 +49,84 @@ extern void R2(void);
 /* "#pragma acc routine" not immediately followed by (a single) function
    declaration or definition.  */
 
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
 int a;
 
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */
 void fn1 (void), fn1b (void);
 
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
 int b, fn2 (void);
 
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
 int b_, fn2_ (void), B_;
 
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */
 int fn3 (void), b2;
 
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
 typedef struct c c;
 
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
 struct d {} d;
 
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */
 void fn1_2 (void), fn1b_2 (void);
 
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
 #pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
 int b_2, fn2_2 (void);
 
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
 int b_2_, fn2_2_ (void), B_2_;
 
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */
 int fn3_2 (void), b2_2;
 
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
 typedef struct c_2 c_2;
 
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
 struct d_2 {} d_2;
 
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
-#pragma acc routine
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine seq
 int fn4 (void);
 
 int fn5a (void);
 int fn5b (void);
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
-#pragma acc routine (fn5a)
-#pragma acc routine (fn5b)
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine (fn5a) seq
+#pragma acc routine (fn5b) seq
 int fn5 (void);
 
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
-#pragma acc routine (fn6a) /* { dg-error ".fn6a. has not been declared" } */
-#pragma acc routine (fn6b) /* { dg-error ".fn6b. has not been declared" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine (fn6a) seq /* { dg-error ".fn6a. has not been declared" } */
+#pragma acc routine (fn6b) seq /* { dg-error ".fn6b. has not been declared" } */
 int fn6 (void);
 
 #ifdef __cplusplus
 
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" "" { target c++ } } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" "" { target c++ } } */
 namespace f {}
 
 namespace g {}
 
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" "" { target c++ } } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" "" { target c++ } } */
 using namespace g;
 
-#pragma acc routine (g) /* { dg-error ".g. does not refer to a function" "" { target c++ } } */
+#pragma acc routine (g) seq /* { dg-error ".g. does not refer to a function" "" { target c++ } } */
 
 #endif /* __cplusplus */
 
-#pragma acc routine (a) /* { dg-error ".a. does not refer to a function" } */
+#pragma acc routine (a) seq /* { dg-error ".a. does not refer to a function" } */
   
-#pragma acc routine (c) /* { dg-error ".c. does not refer to a function" } */
+#pragma acc routine (c) seq /* { dg-error ".c. does not refer to a function" } */
 
 
 /* Static assert.  */
@@ -143,24 +143,24 @@ static_assert(0, ""); /* { dg-error "static assertion failed" "" { target c++11
 #endif
 void f_static_assert();
 /* Check that we already recognized "f_static_assert" as an OpenACC routine.  */
-#pragma acc routine (f_static_assert) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*f_static_assert" "TODO" { xfail *-*-* } } */
+#pragma acc routine (f_static_assert) seq /* { dg-error ".#pragma acc routine. already applied to .\[void \]*f_static_assert" "TODO" { xfail *-*-* } } */
 
 
 /* __extension__ usage.  */
 
-#pragma acc routine
+#pragma acc routine seq
 __extension__ extern void ex1();
 #pragma acc routine (ex1) worker /* { dg-error "has already been marked as an accelerator routine" } */
 
-#pragma acc routine
+#pragma acc routine seq
 __extension__ __extension__ __extension__ __extension__ __extension__ 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" } */
+#pragma acc routine seq /* { 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 (ex3) seq /* { dg-error ".ex3. does not refer to a function" } */
 
 
 /* "#pragma acc routine" must be applied before.  */
@@ -172,11 +172,11 @@ void Foo ()
   Bar ();
 }
 
-#pragma acc routine (Bar) // { dg-error ".#pragma acc routine. must be applied before use" }
+#pragma acc routine (Bar) seq // { dg-error ".#pragma acc routine. must be applied before use" }
 
 #pragma acc routine (Foo) gang // { dg-error ".#pragma acc routine. must be applied before definition" }
 
-#pragma acc routine (Baz) // { dg-error "not been declared" }
+#pragma acc routine (Baz) seq // { dg-error "not been declared" }
 
 
 /* OpenACC declare.  */
@@ -185,7 +185,7 @@ int vb1;		/* { dg-error "directive for use" } */
 extern int vb2;		/* { dg-error "directive for use" } */
 static int vb3;		/* { dg-error "directive for use" } */
 
-#pragma acc routine
+#pragma acc routine seq
 int
 func1 (int a)
 {
@@ -196,7 +196,7 @@ func1 (int a)
   return vb3;
 }
 
-#pragma acc routine
+#pragma acc routine seq
 int
 func2 (int a)
 {
@@ -214,7 +214,7 @@ extern int vb6;			/* { dg-error "clause used in" } */
 static int vb7;			/* { dg-error "clause used in" } */
 #pragma acc declare link (vb7)
 
-#pragma acc routine
+#pragma acc routine seq
 int
 func3 (int a)
 {
@@ -231,7 +231,7 @@ extern int vb9;
 static int vb10;
 #pragma acc declare create (vb10)
 
-#pragma acc routine
+#pragma acc routine seq
 int
 func4 (int a)
 {
@@ -249,7 +249,7 @@ extern int vb12;
 extern int vb13;
 #pragma acc declare device_resident (vb13)
 
-#pragma acc routine
+#pragma acc routine seq
 int
 func5 (int a)
 {
@@ -260,7 +260,7 @@ func5 (int a)
   return vb13;
 }
 
-#pragma acc routine
+#pragma acc routine seq
 int
 func6 (int a)
 {
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
index 8f45499..e32d89d 100644
--- 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
@@ -428,20 +428,20 @@ extern void s_7 (void);
 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 (g_8) seq /* { 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 (w_8) seq /* { 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" } */
+#pragma acc routine (v_8) seq /* { 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) /* { dg-warning "expecting one of" } */
 #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) \
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
index 1803997..7f26ef1 100644
--- 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
@@ -28,46 +28,46 @@ extern void v_1 (void);
 
 #pragma acc routine seq
 extern void s_1_1 (void);
-#pragma acc routine (s_1_1)
+#pragma acc routine (s_1_1) /* { dg-warning "expecting one of" } */
 #pragma acc routine (s_1_1) seq
-#pragma acc routine (s_1_1)
+#pragma acc routine (s_1_1) /* { dg-warning "expecting one of" } */
 #pragma acc routine (s_1_1) seq
 
-#pragma acc routine
+#pragma acc routine  /* { dg-warning "expecting one of" } */
 extern void s_1_2 (void);
-#pragma acc routine (s_1_2)
+#pragma acc routine (s_1_2) /* { dg-warning "expecting one of" } */
 #pragma acc routine (s_1_2) seq
-#pragma acc routine (s_1_2)
+#pragma acc routine (s_1_2) /* { dg-warning "expecting one of" } */
 #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) /* { dg-warning "expecting one of" } */
 #pragma acc routine (s_2_1) seq
-#pragma acc routine (s_2_1)
+#pragma acc routine (s_2_1) /* { dg-warning "expecting one of" } */
 #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) /* { dg-warning "expecting one of" } */
+#pragma acc routine (s_2_2) /* { dg-warning "expecting one of" } */
 #pragma acc routine (s_2_2) seq
-#pragma acc routine (s_2_2)
+#pragma acc routine (s_2_2) /* { dg-warning "expecting one of" } */
 #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) /* { dg-warning "expecting one of" } */
 #pragma acc routine (s_3_1) seq
-#pragma acc routine (s_3_1)
+#pragma acc routine (s_3_1) /* { dg-warning "expecting one of" } */
 #pragma acc routine (s_3_1) seq
 
-#pragma acc routine
+#pragma acc routine seq
 void s_3_2 (void)
 {
 }
-#pragma acc routine (s_3_2)
+#pragma acc routine (s_3_2) /* { dg-warning "expecting one of" } */
 #pragma acc routine (s_3_2) seq
-#pragma acc routine (s_3_2)
+#pragma acc routine (s_3_2) /* { dg-warning "expecting one of" } */
 #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
index 17d6b03..800bcf6 100644
--- a/gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c
@@ -3,46 +3,46 @@
 
 /* { dg-additional-options "-fdump-tree-oaccdevlow" } */
 
-#pragma acc routine nohost
+#pragma acc routine nohost seq
 int THREE(void)
 {
   return 3;
 }
 
-#pragma acc routine (THREE) nohost
+#pragma acc routine (THREE) nohost seq
 
-#pragma acc routine nohost
+#pragma acc routine nohost seq
 extern int THREE(void);
 
 
-#pragma acc routine nohost
+#pragma acc routine nohost seq
 extern void NOTHING(void);
 
-#pragma acc routine (NOTHING) nohost
+#pragma acc routine (NOTHING) nohost seq
 
 void NOTHING(void)
 {
 }
 
-#pragma acc routine nohost
+#pragma acc routine nohost seq
 extern void NOTHING(void);
 
-#pragma acc routine (NOTHING) nohost
+#pragma acc routine (NOTHING) nohost seq
 
 
 extern float ADD(float, float);
 
-#pragma acc routine (ADD) nohost
+#pragma acc routine (ADD) nohost seq
 
 float ADD(float x, float y)
 {
   return x + y;
 }
 
-#pragma acc routine nohost
+#pragma acc routine nohost seq
 extern float ADD(float, float);
 
-#pragma acc routine (ADD) nohost
+#pragma acc routine (ADD) nohost seq
 
 
 /* { 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
index 7402bfc..f172c38 100644
--- a/gcc/testsuite/c-c++-common/goacc/routine-nohost-2.c
+++ b/gcc/testsuite/c-c++-common/goacc/routine-nohost-2.c
@@ -1,41 +1,41 @@
 /* Test invalid usage of the nohost clause for OpenACC routine directive.
    Exercising different variants for declaring routines.  */
 
-#pragma acc routine
+#pragma acc routine seq
 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" } */
+  nohost seq /* { 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" } */
+  nohost seq /* { 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
+#pragma acc routine seq
 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" } */
+  nohost seq /* { 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" } */
+  nohost seq /* { 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" } */
+  nohost seq /* { 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)
+#pragma acc routine (ADD_1) seq
 
 float ADD_1(float x, float y)
 {
@@ -43,55 +43,55 @@ float ADD_1(float x, float 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" } */
+  nohost seq /* { 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" } */
+  nohost seq /* { 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
+  nohost seq
 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 (THREE_2) seq /* { 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" } */
+#pragma acc routine seq /* { 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
+  nohost seq
 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" } */
+#pragma acc routine (NOTHING_2) seq /* { 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" } */
+#pragma acc routine seq /* { 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" } */
+#pragma acc routine (NOTHING_2) seq /* { 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
+  nohost seq
 
 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" } */
+#pragma acc routine seq /* { 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" } */
+#pragma acc routine (ADD_2) seq /* { 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/g++.dg/goacc/routine-1.C b/gcc/testsuite/g++.dg/goacc/routine-1.C
index a73a73d..9257324 100644
--- a/gcc/testsuite/g++.dg/goacc/routine-1.C
+++ b/gcc/testsuite/g++.dg/goacc/routine-1.C
@@ -4,10 +4,10 @@ namespace N
 {
   extern void foo1();
   extern void foo2();
-#pragma acc routine (foo1)
-#pragma acc routine
+#pragma acc routine (foo1) seq
+#pragma acc routine seq
   void foo3()
   {
   }
 }
-#pragma acc routine (N::foo2)
+#pragma acc routine (N::foo2) seq
diff --git a/gcc/testsuite/g++.dg/goacc/routine-2.C b/gcc/testsuite/g++.dg/goacc/routine-2.C
index 939dfba..846f388 100644
--- a/gcc/testsuite/g++.dg/goacc/routine-2.C
+++ b/gcc/testsuite/g++.dg/goacc/routine-2.C
@@ -2,7 +2,7 @@
 
 template <typename T>
 extern T one_d();
-#pragma acc routine (one_d) nohost /* { dg-error "names a set of overloads" } */
+#pragma acc routine (one_d) nohost seq /* { dg-error "names a set of overloads" } */
 
 template <typename T>
 T
@@ -10,41 +10,41 @@ one()
 {
   return 1;
 }
-#pragma acc routine (one) bind(one_d) /* { dg-error "names a set of overloads" } */
+#pragma acc routine (one) bind(one_d) seq /* { dg-error "names a set of overloads" } */
 
 int incr (int);
 float incr (float);
 
-#pragma acc routine (incr) /* { dg-error "names a set of overloads" } */
+#pragma acc routine (incr) seq /* { dg-error "names a set of overloads" } */
 
 
 int sum (int, int);
 
 namespace foo {
-#pragma acc routine (sum)
+#pragma acc routine (sum) seq
   int sub (int, int);
 }
 
-#pragma acc routine (foo::sub)
+#pragma acc routine (foo::sub) seq
 
 /* It's strange to apply a routine directive to subset of overloaded
    functions, but that is permissible in OpenACC 2.x.  */
 
 int decr (int a);
 
-#pragma acc routine
+#pragma acc routine seq
 float decr (float a);
 
 /* Bind clause.  */
 
-#pragma acc routine
+#pragma acc routine seq
 float
 mult (float a, float b)
 {
   return a * b;
 }
 
-#pragma acc routine bind(mult) /* { dg-error "bind identifier 'mult' is not compatible with function 'broken_mult1'" } */
+#pragma acc routine bind(mult) seq /* { dg-error "bind identifier 'mult' is not compatible with function 'broken_mult1'" } */
 float
 broken_mult1 (int a, int b)
 {
@@ -52,51 +52,51 @@ broken_mult1 (int a, int b)
 }
 
 /* This should result in a link error, but it's valid for a compile test.  */
-#pragma acc routine bind("mult")
+#pragma acc routine bind("mult") seq
 float
 broken_mult2 (float a, float b)
 {
   return a + b;
 }
-#pragma acc routine (broken_mult2) bind("mult")
-#pragma acc routine bind("mult")
+#pragma acc routine (broken_mult2) bind("mult") seq
+#pragma acc routine bind("mult") seq
 extern float broken_mult2 (float, float);
 
-#pragma acc routine bind(sum2) /* { dg-error "'sum2' has not been declared" } */
+#pragma acc routine bind(sum2) seq /* { dg-error "'sum2' has not been declared" } */
 int broken_mult3 (int a, int b);
 
-#pragma acc routine bind(foo::sub)
+#pragma acc routine bind(foo::sub) seq
 int broken_mult4 (int a, int b);
-#pragma acc routine (broken_mult4) bind(foo::sub)
-#pragma acc routine bind(foo::sub)
+#pragma acc routine (broken_mult4) bind(foo::sub) seq
+#pragma acc routine bind(foo::sub) seq
 extern int broken_mult4 (int a, int b);
 
 namespace bar
 {
-#pragma acc routine bind (mult)
+#pragma acc routine bind (mult) seq
   float working_mult (float a, float b)
   {
     return a * b;
   }
-#pragma acc routine (working_mult) bind (mult)
-#pragma acc routine bind (mult)
+#pragma acc routine (working_mult) bind (mult) seq
+#pragma acc routine bind (mult) seq
   float working_mult (float, float);
 }
 
-#pragma acc routine
+#pragma acc routine seq
 int div (int, int);
 
-#pragma acc routine
+#pragma acc routine seq
 float div (float, float);
 
-#pragma acc routine bind (div) /* { dg-error "'div' names a set of overloads" } */
+#pragma acc routine bind (div) seq /* { dg-error "'div' names a set of overloads" } */
 float
 my_div (float a, float b)
 {
   return a / b;
 }
 
-#pragma acc routine bind (other_div) /* { dg-error "'other_div' has not been declared" } */
+#pragma acc routine bind (other_div) seq /* { dg-error "'other_div' has not been declared" } */
 float
 my_div2 (float a, float b)
 {
@@ -105,21 +105,21 @@ my_div2 (float a, float b)
 
 int div_var;
 
-#pragma acc routine bind (div_var) /* { dg-error "'div_var' does not refer to a function" } */
+#pragma acc routine bind (div_var) seq /* { dg-error "'div_var' does not refer to a function" } */
 float my_div3 (float, float);
 
-#pragma acc routine bind (div_var) /* { dg-error "'div_var' does not refer to a function" } */
+#pragma acc routine bind (div_var) seq /* { dg-error "'div_var' does not refer to a function" } */
 float my_div4 (float, float);
 
-#pragma acc routine bind (%) /* { dg-error "expected identifier or character string literal" } */
+#pragma acc routine bind (%) seq /* { dg-error "expected identifier or character string literal" } */
 float my_div5 (float, float);
 
-#pragma acc routine bind ("") /* { dg-error "bind argument must not be an empty string" } */
+#pragma acc routine bind ("") seq /* { dg-error "bind argument must not be an empty string" } */
 float my_div6 (float, float);
 
 struct astruct
 {
-  #pragma acc routine /* { dg-error ".#pragma acc routine. must be at file scope" } */
+  #pragma acc routine seq /* { dg-error ".#pragma acc routine. must be at file scope" } */
   int sum (int a, int b)
   {
     return a + b;
@@ -128,7 +128,7 @@ struct astruct
 
 class aclass
 {
-  #pragma acc routine /* { dg-error ".#pragma acc routine. must be at file scope" } */
+  #pragma acc routine seq /* { dg-error ".#pragma acc routine. must be at file scope" } */
   int sum (int a, int b)
   {
     return a + b;
diff --git a/gcc/testsuite/g++.dg/goacc/template.C b/gcc/testsuite/g++.dg/goacc/template.C
index 056398d..4bc2596 100644
--- a/gcc/testsuite/g++.dg/goacc/template.C
+++ b/gcc/testsuite/g++.dg/goacc/template.C
@@ -1,4 +1,4 @@
-#pragma acc routine
+#pragma acc routine seq
 template <typename T> T
 accDouble(int val)
 {
diff --git a/gcc/testsuite/gfortran.dg/goacc/dtype-1.f95 b/gcc/testsuite/gfortran.dg/goacc/dtype-1.f95
index 5919ae4..f24b60f 100644
--- a/gcc/testsuite/gfortran.dg/goacc/dtype-1.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/dtype-1.f95
@@ -108,63 +108,63 @@ end program dtype
 !! ACC ROUTINE:
 
 subroutine sr1 ()
-  !$acc routine device_type (nvidia) gang
+  !$acc routine seq device_type (nvidia) gang
 end subroutine sr1
 
 subroutine sr2 ()
-  !$acc routine dtype (nvidia) worker
+  !$acc routine seq dtype (nvidia) worker
 end subroutine sr2
 
 subroutine sr3 ()
-  !$acc routine device_type (nvidia) vector
+  !$acc routine seq device_type (nvidia) vector
 end subroutine sr3
 
 subroutine sr4 ()
-  !$acc routine device_type (nvidia) seq
+  !$acc routine seq device_type (nvidia) seq
 end subroutine sr4
 
 subroutine sr5 ()
-  !$acc routine dtype (nvidia) bind (foo)
+  !$acc routine seq dtype (nvidia) bind (foo)
 end subroutine sr5
 
 subroutine sr1a ()
-  !$acc routine device_type (nvidia) gang device_type (*) seq
+  !$acc routine seq device_type (nvidia) gang device_type (*) seq
 end subroutine sr1a
 
 subroutine sr2a ()
-  !$acc routine dtype (nvidia) worker dtype (*) seq
+  !$acc routine seq dtype (nvidia) worker dtype (*) seq
 end subroutine sr2a
 
 subroutine sr3a ()
-  !$acc routine dtype (nvidia) vector device_type (*) seq
+  !$acc routine seq dtype (nvidia) vector device_type (*) seq
 end subroutine sr3a
 
 subroutine sr4a ()
-  !$acc routine device_type (nvidia) seq device_type (*) worker
+  !$acc routine seq device_type (nvidia) seq device_type (*) worker
 end subroutine sr4a
 
 subroutine sr5a ()
-  !$acc routine device_type (nvidia) bind (foo) dtype (*) seq
+  !$acc routine seq device_type (nvidia) bind (foo) dtype (*) seq
 end subroutine sr5a
 
 subroutine sr1b ()
-  !$acc routine dtype (gpu) gang dtype (*) seq
+  !$acc routine seq dtype (gpu) gang dtype (*) seq
 end subroutine sr1b
 
 subroutine sr2b ()
-  !$acc routine dtype (gpu) worker device_type (*) seq
+  !$acc routine seq dtype (gpu) worker device_type (*) seq
 end subroutine sr2b
 
 subroutine sr3b ()
-  !$acc routine device_type (gpu) vector device_type (*) seq
+  !$acc routine seq device_type (gpu) vector device_type (*) seq
 end subroutine sr3b
 
 subroutine sr4b ()
-  !$acc routine device_type (gpu) seq device_type (*) worker
+  !$acc routine seq device_type (gpu) seq device_type (*) worker
 end subroutine sr4b
 
 subroutine sr5b ()
-  !$acc routine dtype (gpu) bind (foo) device_type (*) seq
+  !$acc routine seq dtype (gpu) bind (foo) device_type (*) seq
 end subroutine sr5b
 
 ! { dg-final { scan-tree-dump-times "oacc_parallel device_type\\(nvidia\\) \\\[ async\\(1\\) wait\\(1\\) num_gangs\\(100\\) num_workers\\(100\\) vector_length\\(32\\) \\\]" 1 "omplower" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/pr71704.f90 b/gcc/testsuite/gfortran.dg/goacc/pr71704.f90
index 0235e85..92d0c71 100644
--- a/gcc/testsuite/gfortran.dg/goacc/pr71704.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/pr71704.f90
@@ -2,7 +2,7 @@
 ! { dg-do compile }
 
 real function f1 ()
-!$acc routine (f1)
+!$acc routine (f1) seq
   f1 = 1
 end
 
diff --git a/gcc/testsuite/gfortran.dg/goacc/pr72741-2.f b/gcc/testsuite/gfortran.dg/goacc/pr72741-2.f
index 5865144..0f0f7df 100644
--- a/gcc/testsuite/gfortran.dg/goacc/pr72741-2.f
+++ b/gcc/testsuite/gfortran.dg/goacc/pr72741-2.f
@@ -1,5 +1,5 @@
       SUBROUTINE v_1
-!$ACC ROUTINE
+!$ACC ROUTINE SEQ
 !$ACC ROUTINE ! { dg-error "ACC ROUTINE already applied" }
 !$ACC ROUTINE GANG ! { dg-error "ACC ROUTINE already applied" }
 !$ACC ROUTINE SEQ ! { dg-error "ACC ROUTINE already applied" }
@@ -13,7 +13,7 @@
 !$ACC ROUTINE (g_1) GANG
 !$ACC ROUTINE (g_1) VECTOR ! { dg-error "ACC ROUTINE already applied" }
 !$ACC ROUTINE (g_1) SEQ ! { dg-error "ACC ROUTINE already applied" }
-!$ACC ROUTINE (g_1) ! { dg-error "ACC ROUTINE already applied" }
+!$ACC ROUTINE (g_1) WORKER! { dg-error "ACC ROUTINE already applied" }
 !$ACC ROUTINE (g_1) ! { dg-error "ACC ROUTINE already applied" }
 
       CALL v_1
diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-6.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-6.f90
index ee43935..5c1f652 100644
--- a/gcc/testsuite/gfortran.dg/goacc/routine-6.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-6.f90
@@ -4,7 +4,7 @@ module m
 contains
   subroutine subr5 (x) 
   implicit none
-  !$acc routine (subr5)
+  !$acc routine (subr5) seq
   integer, intent(inout) :: x
   if (x < 1) then
      x = 1
@@ -18,7 +18,7 @@ program main
   implicit none
   interface
     function subr6 (x) 
-    !$acc routine (subr6) ! { dg-error "without list is allowed in interface" }
+    !$acc routine (subr6) seq ! { dg-error "without list is allowed in interface" }
     integer, intent (in) :: x
     integer :: subr6
     end function subr6
@@ -26,13 +26,13 @@ program main
   integer, parameter :: n = 10
   integer :: a(n), i
   external :: subr2
-  !$acc routine (subr2)
+  !$acc routine (subr2) seq
 
   external :: R1, R2
-  !$acc routine (R1 R2 R3) ! { dg-error "Syntax error in \\!\\\$ACC ROUTINE \\( NAME \\) at \\(1\\), expecting .\\). after NAME" }
-  !$acc routine (R1, R2, R3) ! { dg-error "Syntax error in \\!\\\$ACC ROUTINE \\( NAME \\) at \\(1\\), expecting .\\). after NAME" }
-  !$acc routine (R1)
-  !$acc routine (R2)
+  !$acc routine (R1 R2 R3) seq ! { dg-error "Syntax error in \\!\\\$ACC ROUTINE \\( NAME \\) at \\(1\\), expecting .\\). after NAME" }
+  !$acc routine (R1, R2, R3) seq ! { dg-error "Syntax error in \\!\\\$ACC ROUTINE \\( NAME \\) at \\(1\\), expecting .\\). after NAME" }
+  !$acc routine (R1) seq
+  !$acc routine (R2) seq
 
   !$acc parallel
   !$acc loop
@@ -44,7 +44,7 @@ program main
 end program main
 
 subroutine subr1 (x) 
-  !$acc routine
+  !$acc routine seq
   integer, intent(inout) :: x
   if (x < 1) then
      x = 1
@@ -63,7 +63,7 @@ subroutine subr2 (x)
 end subroutine subr2
 
 subroutine subr3 (x) 
-  !$acc routine (subr3)
+  !$acc routine (subr3) seq
   integer, intent(inout) :: x
   if (x < 1) then
      x = 1
@@ -73,7 +73,7 @@ subroutine subr3 (x)
 end subroutine subr3
 
 subroutine subr4 (x) 
-  !$acc routine (subr4)
+  !$acc routine (subr4) seq
   integer, intent(inout) :: x
   if (x < 1) then
      x = 1
diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-8.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-8.f90
index c903915..beca43f 100644
--- a/gcc/testsuite/gfortran.dg/goacc/routine-8.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-8.f90
@@ -4,7 +4,7 @@ program main
   interface
      function s_1 (a)
        integer a
-       !$acc routine
+       !$acc routine seq
      end function s_1
   end interface
 
@@ -18,7 +18,7 @@ program main
   interface
      function s_3 (a)
        integer a
-       !$acc routine (s_3) ! { dg-error "Only the ..ACC ROUTINE form without list is allowed in interface block" }
+       !$acc routine (s_3) seq ! { dg-error "Only the ..ACC ROUTINE form without list is allowed in interface block" }
      end function s_3
   end interface
 
diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-9.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-9.f90
index 590e594..4027471 100644
--- a/gcc/testsuite/gfortran.dg/goacc/routine-9.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-9.f90
@@ -7,9 +7,9 @@ contains
   subroutine subr5 (x)
     implicit none
     integer extfunc
-    !$acc routine (subr5)
-    !$acc routine (extfunc)
-    !$acc routine (m1int) ! { dg-error "invalid function name" }
+    !$acc routine (subr5) seq
+    !$acc routine (extfunc) seq
+    !$acc routine (m1int) seq ! { dg-error "invalid function name" }
     integer, intent(inout) :: x
     if (x < 1) then
        x = 1
@@ -29,13 +29,13 @@ program main
   end interface
   integer, parameter :: n = 10
   integer :: a(n), i
-  !$acc routine (subr1) ! { dg-error "invalid function name" }
+  !$acc routine (subr1) seq ! { dg-error "invalid function name" }
   external :: subr2
-  !$acc routine (subr2)
+  !$acc routine (subr2) seq
 
   external :: R1, R2
-  !$acc routine (R1)
-  !$acc routine (R2)
+  !$acc routine (R1) seq
+  !$acc routine (R2) seq
 
   !$acc parallel
   !$acc loop
@@ -47,7 +47,7 @@ program main
 end program main
 
 subroutine subr1 (x)
-  !$acc routine
+  !$acc routine seq
   integer, intent(inout) :: x
   if (x < 1) then
      x = 1
@@ -66,7 +66,7 @@ subroutine subr2 (x)
 end subroutine subr2
 
 subroutine subr3 (x)
-  !$acc routine (subr3)
+  !$acc routine (subr3) seq
   integer, intent(inout) :: x
   if (x < 1) then
      x = 1
@@ -76,7 +76,7 @@ subroutine subr3 (x)
 end subroutine subr3
 
 subroutine subr4 (x)
-  !$acc routine (subr4)
+  !$acc routine (subr4) seq
   integer, intent(inout) :: x
   if (x < 1) then
      x = 1
@@ -86,7 +86,7 @@ subroutine subr4 (x)
 end subroutine subr4
 
 subroutine subr10 (x)
-  !$acc routine (subr10) device ! { dg-error "Unclassifiable OpenACC directive" }
+  !$acc routine (subr10) seq 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-level-of-parallelism-1.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-level-of-parallelism-1.f90
index 364a058..7c245ce 100644
--- a/gcc/testsuite/gfortran.dg/goacc/routine-level-of-parallelism-1.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-level-of-parallelism-1.f90
@@ -9,7 +9,7 @@ subroutine g_1 ! { dg-warning "region is gang partitioned but does not contain g
 end subroutine g_1
 
 subroutine s_1_2a
-  !$acc routine
+  !$acc routine seq
 end subroutine s_1_2a
 
 subroutine s_1_2b
@@ -17,7 +17,7 @@ subroutine s_1_2b
 end subroutine s_1_2b
 
 subroutine s_1_2c
-  !$acc routine (s_1_2c)
+  !$acc routine (s_1_2c) seq
 end subroutine s_1_2c
 
 subroutine s_1_2d
@@ -27,7 +27,7 @@ end subroutine s_1_2d
 module s_2
 contains
   subroutine s_2_1a
-    !$acc routine
+    !$acc routine seq
   end subroutine s_2_1a
 
   subroutine s_2_1b
@@ -35,7 +35,7 @@ contains
   end subroutine s_2_1b
 
   subroutine s_2_1c
-    !$acc routine (s_2_1c)
+    !$acc routine (s_2_1c) seq
   end subroutine s_2_1c
 
   subroutine s_2_1d
@@ -50,7 +50,7 @@ subroutine test
   interface
      function s_3_1a (a)
        integer a
-       !$acc routine
+       !$acc routine seq
      end function s_3_1a
   end interface
 
diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-without-clauses.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-without-clauses.f90
new file mode 100644
index 0000000..570c8ed
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-without-clauses.f90
@@ -0,0 +1,36 @@
+! Test the OpenACC routine directive when it has no gang, worker
+! vector, or seq partitioning clauses.
+
+! { dg-do compile }
+
+subroutine s1
+  !$acc routine ! { dg-warning "Expected one of" }
+end subroutine s1
+
+integer function f1 ()
+  !$acc routine ! { dg-warning "Expected one of" }
+end function f1
+
+module m
+contains
+  subroutine s2
+    !$acc routine ! { dg-warning "Expected one of" }
+  end subroutine s2
+
+  integer function f2 ()
+    !$acc routine ! { dg-warning "Expected one of" }
+  end function f2
+end module m
+
+program t
+  implicit none
+  interface
+     subroutine s3
+       !$acc routine ! { dg-warning "Expected one of" }
+     end subroutine s3
+
+     integer function f3 ()
+       !$acc routine ! { dg-warning "Expected one of" }
+     end function f3
+  end interface
+end program t
diff --git a/libgomp/testsuite/libgomp.oacc-c++/pr71959-a.C b/libgomp/testsuite/libgomp.oacc-c++/pr71959-a.C
index 9486512..82d6c22 100644
--- a/libgomp/testsuite/libgomp.oacc-c++/pr71959-a.C
+++ b/libgomp/testsuite/libgomp.oacc-c++/pr71959-a.C
@@ -8,13 +8,13 @@ struct Iter
   int *point () const asm("_ZNK4Iter5pointEv");
 };
 
-#pragma acc routine
+#pragma acc routine seq
 void  Iter::ctor (int *cursor_)
 {
   cursor = cursor_;
 }
 
-#pragma acc routine
+#pragma acc routine seq
 int *Iter::point () const
 {
   return cursor;
@@ -22,7 +22,7 @@ int *Iter::point () const
 
 void apply (int (*fn)(), Iter out) asm ("_ZN5Apply5applyEPFivE4Iter");
 
-#pragma acc routine
+#pragma acc routine seq
 void apply (int (*fn)(), struct Iter out)
 { *out.point() = fn (); }
 
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-2.c
index 2078a33..4f06ffd 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-2.c
@@ -7,7 +7,7 @@
 float c[N];
 #pragma acc declare device_resident (c)
 
-#pragma acc routine
+#pragma acc routine seq
 float
 subr2 (float a)
 {
@@ -25,7 +25,7 @@ subr2 (float a)
 float b[N];
 #pragma acc declare copyin (b)
 
-#pragma acc routine
+#pragma acc routine seq
 float
 subr1 (float a)
 {
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-3.c
index c3a2187..1b72bda 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-3.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-3.c
@@ -6,7 +6,7 @@
 float *b;
 #pragma acc declare deviceptr (b)
 
-#pragma acc routine
+#pragma acc routine seq
 float *
 subr2 (void)
 {
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-4.c
index 36bf0eb..385b710 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-4.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-4.c
@@ -6,7 +6,7 @@
 float b;
 #pragma acc declare create (b)
 
-#pragma acc routine
+#pragma acc routine seq
 int
 func (int a)
 {
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c
index fe843ec..9cb4665 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c
@@ -8,7 +8,7 @@
 #include <cuda_runtime_api.h>
 #include <cublas_v2.h>
 
-#pragma acc routine
+#pragma acc routine seq
 void
 saxpy (int n, float a, float *x, float *y)
 {
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-default.h b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-default.h
index 55de04b..36e8497 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-default.h
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-default.h
@@ -3,7 +3,7 @@
 #include <string.h>
 #include <stdio.h>
 
-#pragma acc routine
+#pragma acc routine seq
 static int __attribute__ ((noinline)) coord ()
 {
   int res = 0;
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/mode-transitions.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/mode-transitions.c
index f62daf0..d3a5858 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/mode-transitions.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/mode-transitions.c
@@ -1122,7 +1122,7 @@ void t27()
 
 /* Check if worker-single variables get broadcastd to vectors.  */
 
-#pragma acc routine
+#pragma acc routine seq
 float t28_routine ()
 {
   return 2.71;
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-loop-2.h b/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-loop-2.h
index 5691b7e..1934a2b 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-loop-2.h
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-loop-2.h
@@ -2,7 +2,7 @@
 #define VARS
 int a[1500];
 float b[10][15][10];
-#pragma acc routine
+#pragma acc routine seq
 __attribute__((noreturn)) void
 noreturn (void)
 {
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-1.c
index 0f70e26..e45e930 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-1.c
@@ -8,7 +8,7 @@
 
 #include <stdlib.h>
 
-#pragma acc routine
+#pragma acc routine seq
 TEMPLATE
 RETURN_1 fact(TYPE n) RETURN_2
 {
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-3.c
index 2cdd6bf..07b9551 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-3.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-3.c
@@ -6,7 +6,7 @@
 
 #include <stdlib.h>
 
-#pragma acc routine nohost
+#pragma acc routine nohost seq
 int
 foo (int n)
 {
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
index b991bb1..ec4a4a3 100644
--- 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
@@ -13,7 +13,7 @@
 /* "MINUS_TWO" is the device variant for function "TWO".  Similar for "THREE",
    and "FOUR".  Exercising different variants for declaring routines.  */
 
-#pragma acc routine nohost
+#pragma acc routine nohost seq
 extern int MINUS_TWO(void);
 
 int MINUS_TWO(void)
@@ -24,7 +24,7 @@ int MINUS_TWO(void)
 }
 
 extern int TWO(void);
-#pragma acc routine (TWO) bind(MINUS_TWO)
+#pragma acc routine (TWO) bind(MINUS_TWO) seq
 
 int TWO(void)
 {
@@ -34,7 +34,7 @@ int TWO(void)
 }
 
 
-#pragma acc routine nohost
+#pragma acc routine nohost seq
 int MINUS_THREE(void)
 {
   if (!acc_on_device(acc_device_not_host))
@@ -42,7 +42,7 @@ int MINUS_THREE(void)
   return -3;
 }
 
-#pragma acc routine bind(MINUS_THREE)
+#pragma acc routine bind(MINUS_THREE) seq
 extern int THREE(void);
 
 int THREE(void)
@@ -55,7 +55,7 @@ int THREE(void)
 
 /* Due to using a string in the bind clause, we don't need "MINUS_FOUR" in
    scope here.  */
-#pragma acc routine bind("MINUS_FOUR")
+#pragma acc routine bind("MINUS_FOUR") seq
 int FOUR(void)
 {
   if (acc_on_device(acc_device_not_host))
@@ -64,7 +64,7 @@ int FOUR(void)
 }
 
 extern int MINUS_FOUR(void);
-#pragma acc routine (MINUS_FOUR) nohost
+#pragma acc routine (MINUS_FOUR) nohost seq
 
 int MINUS_FOUR(void)
 {
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/host_data-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/host_data-2.f90
index ff09218..4cdaa0c 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/host_data-2.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/host_data-2.f90
@@ -81,7 +81,7 @@ subroutine saxpy (nn, aa, xx, yy)
   integer :: nn
   real*4 :: aa, xx(nn), yy(nn)
   integer i
-  !$acc routine
+  !$acc routine seq
 
   do i = 1, nn
     yy(i) = yy(i) + aa * xx(i)
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/host_data-3.f b/libgomp/testsuite/libgomp.oacc-fortran/host_data-3.f
index 05ed949..fe0b904 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/host_data-3.f
+++ b/libgomp/testsuite/libgomp.oacc-fortran/host_data-3.f
@@ -67,7 +67,7 @@
       integer :: nn
       real*4 :: aa, xx(nn), yy(nn)
       integer i
-!$acc routine
+!$acc routine seq
 
       do i = 1, nn
          yy(i) = yy(i) + aa * xx(i)
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/host_data-4.f90 b/libgomp/testsuite/libgomp.oacc-fortran/host_data-4.f90
index 6e379b5..e192c59 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/host_data-4.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/host_data-4.f90
@@ -21,7 +21,7 @@ contains
     integer :: nn
     real*4 :: aa, xx(nn), yy(nn)
     integer i
-    !$acc routine
+    !$acc routine seq
 
     do i = 1, nn
        yy(i) = yy(i) + aa * xx(i)
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/routine-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/routine-1.f90
index 3390515..3b85391 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/routine-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/routine-1.f90
@@ -3,7 +3,7 @@
 
   interface
     recursive function fact (x)
-      !$acc routine
+      !$acc routine seq
       integer, intent(in) :: x
       integer :: fact
     end function fact
@@ -21,7 +21,7 @@
   end do
 end
 recursive function fact (x) result (res)
-  !$acc routine
+  !$acc routine seq
   integer, intent(in) :: x
   integer :: res
   if (x < 1) then
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/routine-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/routine-2.f90
index 3d418b6..902ab69 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/routine-2.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/routine-2.f90
@@ -4,7 +4,7 @@
   module m1
     contains
     recursive function fact (x) result (res)
-      !$acc routine
+      !$acc routine seq
       integer, intent(in) :: x
       integer :: res
       if (x < 1) then
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/routine-3.f90 b/libgomp/testsuite/libgomp.oacc-fortran/routine-3.f90
index d233a63..c3f6edd 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/routine-3.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/routine-3.f90
@@ -4,7 +4,7 @@
   integer, parameter :: n = 10
   integer :: a(n), i
   integer, external :: fact
-  !$acc routine (fact)
+  !$acc routine (fact) seq
   !$acc parallel
   !$acc loop
   do i = 1, n
@@ -16,7 +16,7 @@
   end do
 end
 recursive function fact (x) result (res)
-  !$acc routine
+  !$acc routine seq
   integer, intent(in) :: x
   integer :: res
   if (x < 1) then
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/routine-4.f90 b/libgomp/testsuite/libgomp.oacc-fortran/routine-4.f90
index 3e5fb09..7e763b8 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/routine-4.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/routine-4.f90
@@ -17,7 +17,7 @@
   end do
 end
 subroutine incr (x)
-  !$acc routine
+  !$acc routine seq
   integer, intent(inout) :: x
   x = x + 1
 end subroutine incr
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/routine-5.f90 b/libgomp/testsuite/libgomp.oacc-fortran/routine-5.f90
index 956da8e..8c5d03a 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/routine-5.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/routine-5.f90
@@ -15,7 +15,7 @@ program main
 contains
 
     function func (n) result (rc)
-    !$acc routine
+    !$acc routine seq
     integer, intent (in) :: n
     integer :: rc
 
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/routine-9.f90 b/libgomp/testsuite/libgomp.oacc-fortran/routine-9.f90
index 95d1a13..6f05e28 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/routine-9.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/routine-9.f90
@@ -6,7 +6,7 @@ program main
   integer, parameter :: n = 10
   integer :: a(n), i
   integer, external :: fact
-  !$acc routine (fact)
+  !$acc routine (fact) seq
   !$acc parallel
   !$acc loop
   do i = 1, n
@@ -20,7 +20,7 @@ end program main
 
 recursive function fact (x) result (res)
   implicit none
-  !$acc routine (fact)
+  !$acc routine (fact) seq
   integer, intent(in) :: x
   integer :: res
   if (x < 1) then

Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]