Use oacc_verify_routine_clauses for C/C++

2018-10-02  Thomas Schwinge  <thomas@codesourcery.com>
	    Cesar Philippidis  <cesar@codesourcery.com>

	gcc/
	* 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/
	* c-parser.c (c_parser_oacc_routine): Normalize order of clauses.
	(c_finish_oacc_routine): Call oacc_verify_routine_clauses.

	gcc/cp/
	* 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/
	* c-c++-common/goacc/routine-level-of-parallelism-1.c: New test.

(cherry picked from openacc-gcc-9-branch commit
1cf459e23c2048b5800cd06e96a6041c17fabbc7)

diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index d1e1ab54ea9e417832a6f6c94831939eec2d276f..d50b791fa963b042c3a7fcb4e3396ba9fc561aaf 100644 (file)
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,11 @@
+2018-10-02  Thomas Schwinge  <thomas@codesourcery.com>
+           Cesar Philippidis  <cesar@codesourcery.com>
+
+       * 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.
+
 2018-10-02  Thomas Schwinge  <thomas@codesourcery.com>
            Cesar Philippidis  <cesar@codesourcery.com>
 

diff --git a/gcc/c/ChangeLog.omp b/gcc/c/ChangeLog.omp
index 023607f1e4d94a8e833d1675fd640fc5349f7c0d..ed7b8b9e73e7c4bc5c9ebebde580a6b178924328 100644 (file)
--- a/gcc/c/ChangeLog.omp
+++ b/gcc/c/ChangeLog.omp
@@ -1,3 +1,9 @@
+2018-10-02  Thomas Schwinge  <thomas@codesourcery.com>
+           Cesar Philippidis  <cesar@codesourcery.com>
+
+       * c-parser.c (c_parser_oacc_routine): Normalize order of clauses.
+       (c_finish_oacc_routine): Call oacc_verify_routine_clauses.
+
 2018-10-02  Thomas Schwinge  <thomas@codesourcery.com>
            Cesar Philippidis  <cesar@codesourcery.com>
 

diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 60e4ed5b63b3ba43544711ed5c9ae90993ee1107..7c2907a765c2cad10d066a6ec95a1329886b990f 100644 (file)
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -15825,6 +15825,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)
        {
@@ -15839,6 +15842,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
@@ -15902,6 +15908,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.omp b/gcc/cp/ChangeLog.omp
index f9ba662d4aa69a71b07fe21d0956176d61583a24..63bae9714ff052c97489039c85c60ccb7f35487a 100644 (file)
--- a/gcc/cp/ChangeLog.omp
+++ b/gcc/cp/ChangeLog.omp
@@ -1,3 +1,10 @@
+2018-10-02  Thomas Schwinge  <thomas@codesourcery.com>
+           Cesar Philippidis  <cesar@codesourcery.com>
+
+       * 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.
+
 2018-10-02  Thomas Schwinge  <thomas@codesourcery.com>
            Cesar Philippidis  <cesar@codesourcery.com>
 

diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 418be26be4b6a1314f9524e786e8f218082bb503..4e3993c0f69f5f4f1f3117919e4811c64d2524a8 100644 (file)
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -40198,6 +40198,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
@@ -40294,6 +40297,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.  */
@@ -40328,6 +40334,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 356772ff4582185d2f410160aa02ec19599d0192..b3cd9cb103d99adda415c92cd4eef886e55bd3ea 100644 (file)
--- a/gcc/omp-general.c
+++ b/gcc/omp-general.c
@@ -608,9 +608,64 @@ 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.
+
+   The chain of clauses returned 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;
+      case OMP_CLAUSE_NOHOST:
+       break;
+      default:
+       gcc_unreachable ();
+      }
+  if (c_level == NULL_TREE)
+    {
+      /* OpenACC 2.5 makes this an error; for the current OpenACC 2.0a
+        implementation add 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's 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 +687,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 60faa5213a27e7166cc53e5b62949f2968f8ae46..4241c33d99e0558bde0c2ccfa015d62a9840652d 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.omp b/gcc/testsuite/ChangeLog.omp
index 2a0d2f4615800213f9df4fcab918c8a50dd0f6fa..adad65edb20fcdba822694eeaef4a10400853667 100644 (file)
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,8 @@
+2018-10-02  Thomas Schwinge  <thomas@codesourcery.com>
+           Cesar Philippidis  <cesar@codesourcery.com>
+
+       * c-c++-common/goacc/routine-level-of-parallelism-1.c: New test.
+
 2018-10-02  Thomas Schwinge  <thomas@codesourcery.com>
            Cesar Philippidis  <cesar@codesourcery.com>
 

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..d71d6e0
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-1.c
@@ -0,0 +1,265 @@
+/* Test various aspects of clauses specifying incompatible levels of
+   parallelism with the OpenACC routine directive.  The Fortran counterpart is
+   ../../gfortran.dg/goacc/routine-level-of-parallelism-1.f.  */
+
+extern void g_1 (void);
+#pragma acc routine (g_1) gang gang /* { dg-error "too many 'gang' clauses" } */
+
+#pragma acc routine worker worker /* { dg-error "too many 'worker' clauses" } */
+void w_1 (void)
+{
+}
+
+#pragma acc routine vector vector /* { dg-error "too many 'vector' clauses" } */
+void v_1 (void)
+{
+}
+
+#pragma acc routine seq seq /* { dg-error "too many 'seq' clauses" } */
+extern void s_1 (void);
+
+
+#pragma acc routine gang gang gang /* { dg-error "too many 'gang' clauses" } */
+void g_2 (void)
+{
+}
+
+#pragma acc routine worker worker worker /* { dg-error "too many 'worker' clauses" } */
+extern void w_2 (void);
+
+extern void v_2 (void);
+#pragma acc routine (v_2) vector vector vector /* { dg-error "too many 'vector' clauses" } */
+
+#pragma acc routine seq seq seq /* { dg-error "too many 'seq' clauses" } */
+void s_2 (void)
+{
+}
+
+
+#pragma acc routine \
+  gang \
+  worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */
+void g_3 (void)
+{
+}
+#pragma acc routine (g_3) /* { 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 } */