[openacc] Repeated use of the OpenACC routine directive

Message ID ebbf5be0-63d7-c048-e595-1de0c6790ae6@codesourcery.com
State New
Headers show
Series
  • [openacc] Repeated use of the OpenACC routine directive
Related show

Commit Message

Cesar Philippidis Oct. 2, 2018, 2:50 p.m.
This is another patch that teaches the C and C++ to emit more errors
involving acc routine clauses. In retrospect, I should have merged it
together with the patch I posted here
<https://gcc.gnu.org/ml/gcc-patches/2018-10/msg00089.html>, however at
the time I thought it would make the patch too large.

Is this patch OK for trunk? I bootstrapped and regtested it for x86_64
Linux with nvptx offloading. This patch is also self-contained to the
OpenACC code path.

Thanks,
Cesar

Patch

[OpenACC] Repeated use of the OpenACC routine directive

2018-XX-YY  Thomas Schwinge  <thomas@codesourcery.com>
	    Cesar Philippidis  <cesar@codesourcery.com>

	gcc/
	* omp-general.h (oacc_verify_routine_clauses): Declare.
	* omp-general.c (oacc_verify_routine_clauses): Change formal
	parameters.  Add checking if already marked as an accelerator
	routine.  Adjust all users.

	gcc/c/
	* c-parser.c (c_finish_oacc_routine): Rework checking if already
	marked as an accelerator routine.

	gcc/cp/
	* parser.c (cp_finalize_oacc_routine): Rework checking if already
	marked as an accelerator routine.

	gcc/testsuite/
	* c-c++-common/goacc/routine-1.c: Update tests.
	* 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: New test.
	* c-c++-common/goacc/routine-nohost-1.c: Update tests.
	* c-c++-common/goacc/routine-nohost-2.c: New test.


(cherry picked from gomp-4_0-branch r239521)

remove bind clause support
---
 gcc/c/c-parser.c                              |  46 ++--
 gcc/cp/parser.c                               |  50 ++--
 gcc/omp-general.c                             | 105 +++++++-
 gcc/omp-general.h                             |   3 +-
 gcc/testsuite/c-c++-common/goacc/routine-1.c  |  10 +-
 gcc/testsuite/c-c++-common/goacc/routine-5.c  |   4 +-
 .../goacc/routine-level-of-parallelism-1.c    | 233 ++++++++++++++++--
 .../goacc/routine-level-of-parallelism-2.c    |  73 ++++++
 .../c-c++-common/goacc/routine-nohost-1.c     |  20 ++
 .../c-c++-common/goacc/routine-nohost-2.c     |  97 ++++++++
 10 files changed, 566 insertions(+), 75 deletions(-)
 create mode 100644 gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-2.c
 create mode 100644 gcc/testsuite/c-c++-common/goacc/routine-nohost-2.c

diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 187a2dec999..3d5cbe76acf 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -15090,35 +15090,39 @@  c_finish_oacc_routine (struct oacc_routine_data *data, tree fndecl,
       return;
     }
 
-  oacc_verify_routine_clauses (&data->clauses, data->loc);
-
-  if (oacc_get_fn_attrib (fndecl))
+  int compatible
+    = oacc_verify_routine_clauses (fndecl, &data->clauses, data->loc,
+				   "#pragma acc routine");
+  if (compatible < 0)
     {
-      error_at (data->loc,
-		"%<#pragma acc routine%> already applied to %qD", fndecl);
       data->error_seen = true;
       return;
     }
-
-  if (TREE_USED (fndecl) || (!is_defn && DECL_SAVED_TREE (fndecl)))
+  if (compatible > 0)
     {
-      error_at (data->loc,
-		TREE_USED (fndecl)
-		? G_("%<#pragma acc routine%> must be applied before use")
-		: G_("%<#pragma acc routine%> must be applied before "
-		     "definition"));
-      data->error_seen = true;
-      return;
     }
+  else
+    {
+      if (TREE_USED (fndecl) || (!is_defn && DECL_SAVED_TREE (fndecl)))
+	{
+	  error_at (data->loc,
+		    TREE_USED (fndecl)
+		    ? G_("%<#pragma acc routine%> must be applied before use")
+		    : G_("%<#pragma acc routine%> must be applied before"
+			 " definition"));
+	  data->error_seen = true;
+	  return;
+	}
 
-  /* Process the routine's dimension clauses.  */
-  tree dims = oacc_build_routine_dims (data->clauses);
-  oacc_replace_fn_attrib (fndecl, dims);
+      /* Set the routine's level of parallelism.  */
+      tree dims = oacc_build_routine_dims (data->clauses);
+      oacc_replace_fn_attrib (fndecl, dims);
 
-  /* Add an "omp declare target" attribute.  */
-  DECL_ATTRIBUTES (fndecl)
-    = tree_cons (get_identifier ("omp declare target"),
-		 data->clauses, DECL_ATTRIBUTES (fndecl));
+      /* Add an "omp declare target" attribute.  */
+      DECL_ATTRIBUTES (fndecl)
+	= tree_cons (get_identifier ("omp declare target"),
+		     data->clauses, DECL_ATTRIBUTES (fndecl));
+    }
 
   /* Remember that we've used this "#pragma acc routine".  */
   data->fndecl_seen = true;
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index d56105ca177..0d314d63cfd 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -38260,36 +38260,42 @@  cp_finalize_oacc_routine (cp_parser *parser, tree fndecl, bool is_defn)
 	  return;
 	}
 
-      oacc_verify_routine_clauses (&parser->oacc_routine->clauses,
-				   parser->oacc_routine->loc);
-
-      if (oacc_get_fn_attrib (fndecl))
+      int compatible
+	= oacc_verify_routine_clauses (fndecl, &parser->oacc_routine->clauses,
+				       parser->oacc_routine->loc,
+				       "#pragma acc routine");
+      if (compatible < 0)
 	{
-	  error_at (parser->oacc_routine->loc,
-		    "%<#pragma acc routine%> already applied to %qD", fndecl);
 	  parser->oacc_routine = NULL;
 	  return;
 	}
-
-      if (TREE_USED (fndecl) || (!is_defn && DECL_SAVED_TREE (fndecl)))
+      if (compatible > 0)
 	{
-	  error_at (parser->oacc_routine->loc,
-		    TREE_USED (fndecl)
-		    ? G_("%<#pragma acc routine%> must be applied before use")
-		    : G_("%<#pragma acc routine%> must be applied before "
-			 "definition"));
-	  parser->oacc_routine = NULL;
-	  return;
 	}
+      else
+	{
+	  if (TREE_USED (fndecl) || (!is_defn && DECL_SAVED_TREE (fndecl)))
+	    {
+	      error_at (parser->oacc_routine->loc,
+			TREE_USED (fndecl)
+			? G_("%<#pragma acc routine%> must be applied before"
+			     " use")
+			: G_("%<#pragma acc routine%> must be applied before"
+			     " definition"));
+	      parser->oacc_routine = NULL;
+	      return;
+	    }
 
-      /* Process the routine's dimension clauses.  */
-      tree dims = oacc_build_routine_dims (parser->oacc_routine->clauses);
-      oacc_replace_fn_attrib (fndecl, dims);
+	  /* Set the routine's level of parallelism.  */
+	  tree dims = oacc_build_routine_dims (parser->oacc_routine->clauses);
+	  oacc_replace_fn_attrib (fndecl, dims);
 
-      /* Add an "omp declare target" attribute.  */
-      DECL_ATTRIBUTES (fndecl)
-	= tree_cons (get_identifier ("omp declare target"),
-		     parser->oacc_routine->clauses, DECL_ATTRIBUTES (fndecl));
+	  /* Add an "omp declare target" attribute.  */
+	  DECL_ATTRIBUTES (fndecl)
+	    = tree_cons (get_identifier ("omp declare target"),
+			 parser->oacc_routine->clauses,
+			 DECL_ATTRIBUTES (fndecl));
+	}
 
       /* Don't unset parser->oacc_routine here: we may still need it to
 	 diagnose wrong usage.  But, remember that we've used this "#pragma acc
diff --git a/gcc/omp-general.c b/gcc/omp-general.c
index 3ea2224957d..5c91ce73a50 100644
--- a/gcc/omp-general.c
+++ b/gcc/omp-general.c
@@ -561,13 +561,17 @@  oacc_set_fn_attrib (tree fn, tree clauses, vec<tree> *args)
 
 /* Verify OpenACC routine clauses.
 
-   The chain of clauses returned will contain exactly one clause specifying the
-   level of parallelism.  */
+   Returns 0 if FNDECL should be marked as an accelerator routine, 1 if it has
+   already been marked in compatible way, and -1 if incompatible.  Upon
+   returning, the chain of clauses will contain exactly one clause specifying
+   the level of parallelism.  */
 
-void
-oacc_verify_routine_clauses (tree *clauses, location_t loc)
+int
+oacc_verify_routine_clauses (tree fndecl, tree *clauses, location_t loc,
+			     const char *routine_str)
 {
   tree c_level = NULL_TREE;
+  tree c_nohost = NULL_TREE;
   tree c_p = NULL_TREE;
   for (tree c = *clauses; c; c_p = c, c = OMP_CLAUSE_CHAIN (c))
     switch (OMP_CLAUSE_CODE (c))
@@ -601,6 +605,8 @@  oacc_verify_routine_clauses (tree *clauses, location_t loc)
 	  }
 	break;
       case OMP_CLAUSE_NOHOST:
+	/* Don't bother with duplicate clauses at this point.  */
+	c_nohost = c;
 	break;
       default:
 	gcc_unreachable ();
@@ -613,6 +619,97 @@  oacc_verify_routine_clauses (tree *clauses, location_t loc)
       OMP_CLAUSE_CHAIN (c_level) = *clauses;
       *clauses = c_level;
     }
+  /* In *clauses, we now have exactly one clause specifying the level of
+     parallelism.  */
+
+  /* Still got some work to do for Fortran...  */
+  if (fndecl == NULL_TREE)
+    return 0;
+
+  tree attr
+    = lookup_attribute ("omp declare target", DECL_ATTRIBUTES (fndecl));
+  if (attr != NULL_TREE)
+    {
+      /* If a "#pragma acc routine" has already been applied, just verify
+	 this one for compatibility.  */
+      /* Collect previous directive's clauses.  */
+      tree c_level_p = NULL_TREE;
+      tree c_nohost_p = NULL_TREE;
+      for (tree c = TREE_VALUE (attr); c; c = OMP_CLAUSE_CHAIN (c))
+	switch (OMP_CLAUSE_CODE (c))
+	  {
+	  case OMP_CLAUSE_GANG:
+	  case OMP_CLAUSE_WORKER:
+	  case OMP_CLAUSE_VECTOR:
+	  case OMP_CLAUSE_SEQ:
+	    gcc_checking_assert (c_level_p == NULL_TREE);
+	    c_level_p = c;
+	    break;
+	  case OMP_CLAUSE_NOHOST:
+	    /* Don't bother with duplicate clauses at this point.  */
+	    c_nohost_p = c;
+	    break;
+	    break;
+	  default:
+	    gcc_unreachable ();
+	  }
+      gcc_checking_assert (c_level_p != NULL_TREE);
+      /* ..., and compare to current directive's, which we've already collected
+	 above.  */
+      tree c_diag;
+      tree c_diag_p;
+      /* Matching level of parallelism?  */
+      if (OMP_CLAUSE_CODE (c_level) != OMP_CLAUSE_CODE (c_level_p))
+	{
+	  c_diag = c_level;
+	  c_diag_p = c_level_p;
+	  goto incompatible;
+	}
+      /* Matching nohost clauses?  */
+      if ((c_nohost == NULL_TREE) != (c_nohost_p == NULL_TREE))
+	{
+	  c_diag = c_nohost;
+	  c_diag_p = c_nohost_p;
+	  goto incompatible;
+	}
+      /* Compatible.  */
+      return 1;
+
+    incompatible:
+      if (c_diag != NULL_TREE)
+	error_at (OMP_CLAUSE_LOCATION (c_diag),
+		  "incompatible %qs clause when applying"
+		  " %<%s%> to %qD, which has already been"
+		  " marked as an accelerator routine",
+		  omp_clause_code_name[OMP_CLAUSE_CODE (c_diag)],
+		  routine_str, fndecl);
+      else if (c_diag_p != NULL_TREE)
+	error_at (loc,
+		  "missing %qs clause when applying"
+		  " %<%s%> to %qD, which has already been"
+		  " marked as an accelerator routine",
+		  omp_clause_code_name[OMP_CLAUSE_CODE (c_diag_p)],
+		  routine_str, fndecl);
+      else
+	gcc_unreachable ();
+      if (c_diag_p != NULL_TREE)
+	inform (OMP_CLAUSE_LOCATION (c_diag_p),
+		"... with %qs clause here",
+		omp_clause_code_name[OMP_CLAUSE_CODE (c_diag_p)]);
+      else
+	{
+	  /* In the front ends, we don't preserve location information for the
+	     OpenACC routine directive itself.  However, that of c_level_p
+	     should be close.  */
+	  location_t loc_routine = OMP_CLAUSE_LOCATION (c_level_p);
+	  inform (loc_routine, "... without %qs clause near to here",
+		  omp_clause_code_name[OMP_CLAUSE_CODE (c_diag)]);
+	}
+      /* Incompatible.  */
+      return -1;
+    }
+
+  return 0;
 }
 
 /*  Process the OpenACC routine's clauses to generate an attribute
diff --git a/gcc/omp-general.h b/gcc/omp-general.h
index e0de288dd9a..1419b66bebb 100644
--- a/gcc/omp-general.h
+++ b/gcc/omp-general.h
@@ -83,7 +83,8 @@  extern int omp_max_simt_vf (void);
 extern tree oacc_launch_pack (unsigned code, tree device, unsigned op);
 extern void oacc_replace_fn_attrib (tree fn, tree dims);
 extern void oacc_set_fn_attrib (tree fn, tree clauses, vec<tree> *args);
-extern void oacc_verify_routine_clauses (tree *, location_t);
+extern int oacc_verify_routine_clauses (tree, tree *, location_t,
+					const char *);
 extern tree oacc_build_routine_dims (tree clauses);
 extern tree oacc_get_fn_attrib (tree fn);
 extern bool offloading_function_p (tree fn);
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-1.c b/gcc/testsuite/c-c++-common/goacc/routine-1.c
index db1322e11ca..23e273c58f4 100644
--- a/gcc/testsuite/c-c++-common/goacc/routine-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/routine-1.c
@@ -20,11 +20,19 @@  void seq (void)
 {
 }
 
-#pragma acc routine nohost
+#pragma acc routine \
+  nohost
 void nohost (void)
 {
 }
 
+#pragma acc routine \
+  nohost
+extern void nohost (void);
+
+#pragma acc routine (nohost) \
+  nohost
+
 int main ()
 {
 #pragma acc kernels num_gangs (32) num_workers (32) vector_length (32)
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-5.c b/gcc/testsuite/c-c++-common/goacc/routine-5.c
index b759db3292d..8b7da7dd453 100644
--- a/gcc/testsuite/c-c++-common/goacc/routine-5.c
+++ b/gcc/testsuite/c-c++-common/goacc/routine-5.c
@@ -150,13 +150,13 @@  void f_static_assert();
 
 #pragma acc routine seq
 __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 seq
 __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 seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
 __extension__ int ex3;
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 d71d6e088f7..d1258750ae0 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
@@ -42,10 +42,10 @@  void s_2 (void)
 void g_3 (void)
 {
 }
-#pragma acc routine (g_3) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*g_3." } */ \
+#pragma acc routine (g_3) \
   gang \
   seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */
-#pragma acc routine (g_3) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*g_3." } */ \
+#pragma acc routine (g_3) \
   gang \
   vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */
 
@@ -53,10 +53,10 @@  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) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*w_3." } */ \
+#pragma acc routine (w_3) \
   worker \
   gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */
-#pragma acc routine (w_3) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*w_3." } */ \
+#pragma acc routine (w_3) \
   worker \
   seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */
 
@@ -66,10 +66,10 @@  extern void w_3 (void);
 void v_3 (void)
 {
 }
-#pragma acc routine (v_3) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*v_3." } */ \
+#pragma acc routine (v_3) \
   vector \
   worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */
-#pragma acc routine (v_3) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*v_3." } */ \
+#pragma acc routine (v_3) \
   vector \
   gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */
 
@@ -77,10 +77,10 @@  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) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*s_3." } */ \
+#pragma acc routine (s_3) \
   seq \
   vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */
-#pragma acc routine (s_3) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*s_3." } */ \
+#pragma acc routine (s_3) \
   seq \
   worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */
 
@@ -91,12 +91,12 @@  extern void s_3 (void);
   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) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*g_4." } */ \
+#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) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*g_4." } */ \
+#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" } */ \
@@ -108,12 +108,12 @@  extern void w_4 (void);
   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) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*w_4." } */ \
+#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) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*w_4." } */ \
+#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" } */ \
@@ -127,12 +127,12 @@  extern void w_4 (void);
 void v_4 (void)
 {
 }
-#pragma acc routine (v_4) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*v_4." } */ \
+#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) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*v_4." } */ \
+#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" } */ \
@@ -146,12 +146,12 @@  void v_4 (void)
 void s_4 (void)
 {
 }
-#pragma acc routine (s_4) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*s_4." } */ \
+#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) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*s_4." } */ \
+#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" } */ \
@@ -169,7 +169,7 @@  void s_4 (void)
 void g_5 (void)
 {
 }
-#pragma acc routine (g_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*g_5." } */ \
+#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 *-*-* } .-1 } */ \
@@ -177,7 +177,7 @@  void g_5 (void)
   /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
   seq seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
   /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
-#pragma acc routine (g_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*g_5." } */ \
+#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 *-*-* } .-1 } */ \
@@ -195,7 +195,7 @@  void g_5 (void)
   gang gang /* { dg-error "too many 'gang' clauses" } */ \
   /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
 extern void w_5 (void);
-#pragma acc routine (w_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*w_5." } */ \
+#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 *-*-* } .-1 } */ \
@@ -203,7 +203,7 @@  extern void w_5 (void);
   /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
   gang gang /* { dg-error "too many 'gang' clauses" } */ \
   /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
-#pragma acc routine (w_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*w_5." } */ \
+#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 *-*-* } .-1 } */ \
@@ -221,7 +221,7 @@  extern void w_5 (void);
   gang gang /* { dg-error "too many 'gang' clauses" } */ \
   /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
 extern void v_5 (void);
-#pragma acc routine (v_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*v_5." } */ \
+#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 *-*-* } .-1 } */ \
@@ -229,7 +229,7 @@  extern void v_5 (void);
   /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
   gang gang /* { dg-error "too many 'gang' clauses" } */ \
   /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
-#pragma acc routine (v_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*v_5." } */ \
+#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 *-*-* } .-1 } */ \
@@ -247,7 +247,7 @@  extern void s_5 (void);
   /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
   gang gang /* { dg-error "too many 'gang' clauses" } */ \
   /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
-#pragma acc routine (s_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*s_5." } */ \
+#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 *-*-* } .-1 } */ \
@@ -255,7 +255,7 @@  extern void s_5 (void);
   /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
   worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
   /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
-#pragma acc routine (s_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*s_5." } */ \
+#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 *-*-* } .-1 } */ \
@@ -263,3 +263,188 @@  extern void s_5 (void);
   /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
   gang gang /* { dg-error "too many 'gang' clauses" } */ \
   /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
+
+
+/* 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 *-*-* } .-1 } */ \
+  vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
+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 *-*-* } .-1 } */ \
+  gang gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  seq seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
+#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 *-*-* } .-1 } */ \
+  gang gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  worker worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
+
+#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 *-*-* } .-1 } */ \
+  vector vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
+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 *-*-* } .-1 } */ \
+  worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
+#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 *-*-* } .-1 } */ \
+  worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
+
+#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 *-*-* } .-1 } */ \
+  seq seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
+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 *-*-* } .-1 } */ \
+  vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
+#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 *-*-* } .-1 } */ \
+  vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
+
+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 *-*-* } .-1 } */ \
+  worker worker worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
+#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 *-*-* } .-1 } */ \
+  seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
+#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 *-*-* } .-1 } */ \
+  seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
+
+
+/* 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 00000000000..36626229299
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-2.c
@@ -0,0 +1,73 @@ 
+/* 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) /* { dg-warning "region is gang partitioned but does not contain gang partitioned code" } */
+/* { dg-bogus "region is worker partitioned but does not contain worker partitioned code" "worker partitioned" { xfail *-*-* } .-1 } */
+/* { dg-bogus "region is vector partitioned but does not contain vector partitioned code" "worker partitioned" { xfail *-*-* } .-2 } */
+{
+}
+#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
index 9baa56cb206..17d6b0387e3 100644
--- a/gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c
@@ -9,13 +9,27 @@  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
@@ -25,4 +39,10 @@  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 00000000000..7402bfc1b41
--- /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" } */
-- 
2.17.1