This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
[patch,openacc] Use oacc_verify_routine_clauses for C/C++
- From: Cesar Philippidis <cesar at codesourcery dot com>
- To: "gcc-patches at gcc dot gnu dot org" <gcc-patches at gcc dot gnu dot org>, Thomas Schwinge <thomas at codesourcery dot com>
- Date: Tue, 2 Oct 2018 07:06:12 -0700
- Subject: [patch,openacc] Use oacc_verify_routine_clauses for C/C++
This patch introduces a new oacc_verify_routine_clauses function that
reports errors if the user abuses the gang, worker and vector clauses
for acc routine directives in C/C++. Fortran is a little different,
because the FE has it's own IR. So, while it would be possible to defer
checking for gang, worker, vector parallelism until a tree node is
created for a function, we'd still have problems of verifying the
parallelism for functions and subroutines defined and declared inside
modules. The C and C++ FE's are similar enough were they can share a
common function.
Is this OK for trunk? I bootstrapped and regression tested it for x86_64
Linux with nvptx offloading. This is only touches the OpenACC code path.
Cesar
[OpenACC] Use oacc_verify_routine_clauses for C/C++
2018-XX-YY 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 gomp-4_0-branch r239520)
---
gcc/c/c-parser.c | 8 +
gcc/cp/parser.c | 9 +
gcc/omp-general.c | 69 ++++-
gcc/omp-general.h | 1 +
.../goacc/routine-level-of-parallelism-1.c | 265 ++++++++++++++++++
5 files changed, 342 insertions(+), 10 deletions(-)
create mode 100644 gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-1.c
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 3ca8fe71cc4..3517cb783d9 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -14999,6 +14999,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)
{
@@ -15013,6 +15016,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
@@ -15076,6 +15082,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/parser.c b/gcc/cp/parser.c
index 241226d8c21..fa7ee7798ae 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -38117,6 +38117,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
@@ -38213,6 +38216,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. */
@@ -38247,6 +38253,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 cabbbbc6de2..3ea2224957d 100644
--- a/gcc/omp-general.c
+++ b/gcc/omp-general.c
@@ -559,9 +559,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
@@ -583,16 +638,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 66f0a33c2e2..e0de288dd9a 100644
--- a/gcc/omp-general.h
+++ b/gcc/omp-general.h
@@ -83,6 +83,7 @@ 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 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-level-of-parallelism-1.c b/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-1.c
new file mode 100644
index 00000000000..d71d6e088f7
--- /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 } */
--
2.17.1