[PR89433] Use 'oacc_verify_routine_clauses' for C/C++ OpenACC 'routine' directives

	gcc/
	PR middle-end/89433
	* omp-general.c (oacc_build_routine_dims): Move some of its
	processing into...
	(oacc_verify_routine_clauses): ... this new function.
	* omp-general.h (oacc_verify_routine_clauses): New prototype.
	gcc/c/
	PR c/89433
	* c-parser.c (c_parser_oacc_routine): Normalize order of clauses.
	(c_finish_oacc_routine): Call oacc_verify_routine_clauses.
	gcc/cp/
	PR c++/89433
	* parser.c (cp_parser_oacc_routine)
	(cp_parser_late_parsing_oacc_routine): Normalize order of clauses.
	(cp_finalize_oacc_routine): Call oacc_verify_routine_clauses.
	gcc/testsuite/
	PR testsuite/89433
	* c-c++-common/goacc/routine-2.c: Update, and move some test
	into...
	* c-c++-common/goacc/routine-level-of-parallelism-1.c: ... this
	new file.

From-SVN: r271344

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index 5fc9724..b8a4358 100644 (file)
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,11 @@
+2019-05-17  Thomas Schwinge  <thomas@codesourcery.com>
+
+       PR middle-end/89433
+       * omp-general.c (oacc_build_routine_dims): Move some of its
+       processing into...
+       (oacc_verify_routine_clauses): ... this new function.
+       * omp-general.h (oacc_verify_routine_clauses): New prototype.
+
 2019-05-17  Iain Sandoe  <iain@sandoe.co.uk>
 
        * config/rs6000/rs6000.c (machopic_output_stub): Adjust the

diff --git a/gcc/c/ChangeLog b/gcc/c/ChangeLog
index f0cab2e..1393e8f 100644 (file)
--- a/gcc/c/ChangeLog
+++ b/gcc/c/ChangeLog
@@ -1,6 +1,10 @@
 2019-05-17  Thomas Schwinge  <thomas@codesourcery.com>
 
        PR c/89433
+       * c-parser.c (c_parser_oacc_routine): Normalize order of clauses.
+       (c_finish_oacc_routine): Call oacc_verify_routine_clauses.
+
+       PR c/89433
        * c-parser.c (c_finish_oacc_routine): Refer to OpenACC 'routine'
        clauses from "omp declare target" attribute.
 

diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 3cbbb19..8337f1c 100644 (file)
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -15801,6 +15801,9 @@ c_parser_oacc_routine (c_parser *parser, enum pragma_context context)
       data.clauses
        = c_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK,
                                     "#pragma acc routine");
+      /* The clauses are in reverse order; fix that to make later diagnostic
+        emission easier.  */
+      data.clauses = nreverse (data.clauses);
 
       if (TREE_CODE (decl) != FUNCTION_DECL)
        {
@@ -15815,6 +15818,9 @@ c_parser_oacc_routine (c_parser *parser, enum pragma_context context)
       data.clauses
        = c_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK,
                                     "#pragma acc routine");
+      /* The clauses are in reverse order; fix that to make later diagnostic
+        emission easier.  */
+      data.clauses = nreverse (data.clauses);
 
       /* Emit a helpful diagnostic if there's another pragma following this
         one.  Also don't allow a static assertion declaration, as in the
@@ -15878,6 +15884,8 @@ 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))
     {
       error_at (data->loc,

diff --git a/gcc/cp/ChangeLog b/gcc/cp/ChangeLog
index 40622ac..2f1e06c 100644 (file)
--- a/gcc/cp/ChangeLog
+++ b/gcc/cp/ChangeLog
@@ -1,6 +1,11 @@
 2019-05-17  Thomas Schwinge  <thomas@codesourcery.com>
 
        PR c++/89433
+       * parser.c (cp_parser_oacc_routine)
+       (cp_parser_late_parsing_oacc_routine): Normalize order of clauses.
+       (cp_finalize_oacc_routine): Call oacc_verify_routine_clauses.
+
+       PR c++/89433
        * parser.c (cp_finalize_oacc_routine): Refer to OpenACC 'routine'
        clauses from "omp declare target" attribute.
 

diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 15424b6..aa6507e 100644 (file)
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -40136,6 +40136,9 @@ cp_parser_oacc_routine (cp_parser *parser, cp_token *pragma_tok,
        = cp_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK,
                                      "#pragma acc routine",
                                      cp_lexer_peek_token (parser->lexer));
+      /* The clauses are in reverse order; fix that to make later diagnostic
+        emission easier.  */
+      data.clauses = nreverse (data.clauses);
 
       if (decl && is_overloaded_fn (decl)
          && (TREE_CODE (decl) != FUNCTION_DECL
@@ -40232,6 +40235,9 @@ cp_parser_late_parsing_oacc_routine (cp_parser *parser, tree attrs)
   parser->oacc_routine->clauses
     = cp_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK,
                                  "#pragma acc routine", pragma_tok);
+  /* The clauses are in reverse order; fix that to make later diagnostic
+     emission easier.  */
+  parser->oacc_routine->clauses = nreverse (parser->oacc_routine->clauses);
   cp_parser_pop_lexer (parser);
   /* Later, cp_finalize_oacc_routine will process the clauses, and then set
      fndecl_seen.  */
@@ -40266,6 +40272,9 @@ 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))
        {
          error_at (parser->oacc_routine->loc,

diff --git a/gcc/omp-general.c b/gcc/omp-general.c
index 356772f..f1d859b 100644 (file)
--- a/gcc/omp-general.c
+++ b/gcc/omp-general.c
@@ -608,9 +608,61 @@ oacc_set_fn_attrib (tree fn, tree clauses, vec<tree> *args)
     }
 }
 
-/*  Process the routine's dimension clauess to generate an attribute
-    value.  Issue diagnostics as appropriate.  We default to SEQ
-    (OpenACC 2.5 clarifies this). All dimensions have a size of zero
+/* Verify OpenACC routine clauses.
+
+   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)
+{
+  tree c_level = 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))
+      {
+      case OMP_CLAUSE_GANG:
+      case OMP_CLAUSE_WORKER:
+      case OMP_CLAUSE_VECTOR:
+      case OMP_CLAUSE_SEQ:
+       if (c_level == NULL_TREE)
+         c_level = c;
+       else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_CODE (c_level))
+         {
+           /* This has already been diagnosed in the front ends.  */
+           /* Drop the duplicate clause.  */
+           gcc_checking_assert (c_p != NULL_TREE);
+           OMP_CLAUSE_CHAIN (c_p) = OMP_CLAUSE_CHAIN (c);
+           c = c_p;
+         }
+       else
+         {
+           error_at (OMP_CLAUSE_LOCATION (c),
+                     "%qs specifies a conflicting level of parallelism",
+                     omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+           inform (OMP_CLAUSE_LOCATION (c_level),
+                   "... to the previous %qs clause here",
+                   omp_clause_code_name[OMP_CLAUSE_CODE (c_level)]);
+           /* Drop the conflicting clause.  */
+           gcc_checking_assert (c_p != NULL_TREE);
+           OMP_CLAUSE_CHAIN (c_p) = OMP_CLAUSE_CHAIN (c);
+           c = c_p;
+         }
+       break;
+      default:
+       gcc_unreachable ();
+      }
+  if (c_level == NULL_TREE)
+    {
+      /* Default to an implicit 'seq' clause.  */
+      c_level = build_omp_clause (loc, OMP_CLAUSE_SEQ);
+      OMP_CLAUSE_CHAIN (c_level) = *clauses;
+      *clauses = c_level;
+    }
+}
+
+/*  Process the OpenACC 'routine' directive clauses to generate an attribute
+    for the level of parallelism.  All dimensions have a size of zero
     (dynamic).  TREE_PURPOSE is set to indicate whether that dimension
     can have a loop partitioned on it.  non-zero indicates
     yes, zero indicates no.  By construction once a non-zero has been
@@ -632,16 +684,10 @@ oacc_build_routine_dims (tree clauses)
     for (ix = GOMP_DIM_MAX + 1; ix--;)
       if (OMP_CLAUSE_CODE (clauses) == ids[ix])
        {
-         if (level >= 0)
-           error_at (OMP_CLAUSE_LOCATION (clauses),
-                     "multiple loop axes specified for routine");
          level = ix;
          break;
        }
-
-  /* Default to SEQ.  */
-  if (level < 0)
-    level = GOMP_DIM_MAX;
+  gcc_checking_assert (level >= 0);
 
   tree dims = NULL_TREE;
 

diff --git a/gcc/omp-general.h b/gcc/omp-general.h
index 60faa52..4241c33 100644 (file)
--- a/gcc/omp-general.h
+++ b/gcc/omp-general.h
@@ -84,6 +84,7 @@ extern tree oacc_launch_pack (unsigned code, tree device, unsigned op);
 extern tree oacc_replace_fn_attrib_attr (tree attribs, tree dims);
 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 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/ChangeLog b/gcc/testsuite/ChangeLog
index 0f44c6d..4b07888 100644 (file)
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,6 +1,12 @@
 2019-05-17  Thomas Schwinge  <thomas@codesourcery.com>
 
        PR testsuite/89433
+       * c-c++-common/goacc/routine-2.c: Update, and move some test
+       into...
+       * c-c++-common/goacc/routine-level-of-parallelism-1.c: ... this
+       new file.
+
+       PR testsuite/89433
        * c-c++-common/goacc/classify-routine.c: Update.
        * gfortran.dg/goacc/classify-routine.f95: Likewise.
 

diff --git a/gcc/testsuite/c-c++-common/goacc/routine-2.c b/gcc/testsuite/c-c++-common/goacc/routine-2.c
index fc5eb11..be1510a 100644 (file)
--- a/gcc/testsuite/c-c++-common/goacc/routine-2.c
+++ b/gcc/testsuite/c-c++-common/goacc/routine-2.c
@@ -1,21 +1,3 @@
-#pragma acc routine gang worker /* { dg-error "multiple loop axes" } */
-void gang (void)
-{
-}
-
-#pragma acc routine worker vector /* { dg-error "multiple loop axes" } */
-void worker (void)
-{
-}
-
-#pragma acc routine vector seq /* { dg-error "multiple loop axes" } */
-void vector (void)
-{
-}
-
-#pragma acc routine seq gang /* { dg-error "multiple loop axes" } */
-void seq (void)
-{
-}
+/* Test invalid use of the OpenACC 'routine' directive.  */
 
 #pragma acc routine (nothing) gang /* { dg-error "not been declared" } */

diff --git a/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-1.c b/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-1.c
new file mode 100644 (file)
index 0000000..ab0414b
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-1.c
@@ -0,0 +1,264 @@
+/* Test various aspects of clauses specifying incompatible levels of
+   parallelism with the OpenACC 'routine' directive.  */
+
+extern void g_1 (void);
+#pragma acc routine (g_1) gang gang /* { dg-error "too many 'gang' clauses" } */
+
+#pragma acc routine worker worker /* { dg-error "too many 'worker' clauses" } */
+void w_1 (void)
+{
+}
+
+#pragma acc routine vector vector /* { dg-error "too many 'vector' clauses" } */
+void v_1 (void)
+{
+}
+
+#pragma acc routine seq seq /* { dg-error "too many 'seq' clauses" } */
+extern void s_1 (void);
+
+
+#pragma acc routine gang gang gang /* { dg-error "too many 'gang' clauses" } */
+void g_2 (void)
+{
+}
+
+#pragma acc routine worker worker worker /* { dg-error "too many 'worker' clauses" } */
+extern void w_2 (void);
+
+extern void v_2 (void);
+#pragma acc routine (v_2) vector vector vector /* { dg-error "too many 'vector' clauses" } */
+
+#pragma acc routine seq seq seq /* { dg-error "too many 'seq' clauses" } */
+void s_2 (void)
+{
+}
+
+
+#pragma acc routine \
+  gang \
+  worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */
+void g_3 (void)
+{
+}
+#pragma acc routine (g_3) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*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." } */ \
+  gang \
+  vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */
+
+extern void w_3 (void);
+#pragma acc routine (w_3) \
+  worker \
+  vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */
+#pragma acc routine (w_3) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*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." } */ \
+  worker \
+  seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */
+
+#pragma acc routine \
+  vector \
+  seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */
+void v_3 (void)
+{
+}
+#pragma acc routine (v_3) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*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." } */ \
+  vector \
+  gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */
+
+extern void s_3 (void);
+#pragma acc routine (s_3) \
+  seq \
+  gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */
+#pragma acc routine (s_3) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*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." } */ \
+  seq \
+  worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */
+
+
+#pragma acc routine \
+  gang gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */ \
+  vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ \
+  seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */
+extern void g_4 (void);
+#pragma acc routine (g_4) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*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." } */ \
+  gang gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */ \
+  vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ \
+  worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */
+
+extern void w_4 (void);
+#pragma acc routine (w_4) \
+  worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */ \
+  vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ \
+  gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */
+#pragma acc routine (w_4) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*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." } */ \
+  worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */ \
+  gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */ \
+  vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */
+
+#pragma acc routine \
+  vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */ \
+  seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */ \
+  gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */
+void v_4 (void)
+{
+}
+#pragma acc routine (v_4) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*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." } */ \
+  vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */ \
+  seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */ \
+  worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */
+
+#pragma acc routine \
+  seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ \
+  worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */ \
+  gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */
+void s_4 (void)
+{
+}
+#pragma acc routine (s_4) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*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." } */ \
+  seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */ \
+  vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ \
+  gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */
+
+
+#pragma acc routine \
+  gang gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-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_5 (void)
+{
+}
+#pragma acc routine (g_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*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 } */ \
+  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_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*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 } */ \
+  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_5 (void);
+#pragma acc routine (w_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*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 } */ \
+  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_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*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 } */ \
+  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_5 (void);
+#pragma acc routine (v_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*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 } */ \
+  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_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*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 } */ \
+  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_5 (void);
+#pragma acc routine (s_5) \
+  seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } .-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_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*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 } */ \
+  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_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*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 } */ \
+  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 } */