From: Thomas Schwinge Date: Fri, 17 May 2019 19:13:15 +0000 (+0200) Subject: [PR89433] Use 'oacc_verify_routine_clauses' for C/C++ OpenACC 'routine' directives X-Git-Url: https://git.libre-soc.org/?a=commitdiff_plain;h=5bf04509f437ff175c001a1c84a13b3a845174eb;p=gcc.git [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 5fc97248662..b8a43586052 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,11 @@ +2019-05-17 Thomas Schwinge + + 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 * config/rs6000/rs6000.c (machopic_output_stub): Adjust the diff --git a/gcc/c/ChangeLog b/gcc/c/ChangeLog index f0cab2e6592..1393e8f47fd 100644 --- a/gcc/c/ChangeLog +++ b/gcc/c/ChangeLog @@ -1,5 +1,9 @@ 2019-05-17 Thomas Schwinge + 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 3cbbb199bdd..8337f1cce0c 100644 --- 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 40622acc0ff..2f1e06ca458 100644 --- a/gcc/cp/ChangeLog +++ b/gcc/cp/ChangeLog @@ -1,5 +1,10 @@ 2019-05-17 Thomas Schwinge + 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 15424b6b633..aa6507e42f4 100644 --- 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 356772ff458..f1d859b0275 100644 --- a/gcc/omp-general.c +++ b/gcc/omp-general.c @@ -608,9 +608,61 @@ oacc_set_fn_attrib (tree fn, tree clauses, vec *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 60faa5213a2..4241c33d99e 100644 --- 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 *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 0f44c6d4ab0..4b07888dadb 100644 --- a/gcc/testsuite/ChangeLog +++ b/gcc/testsuite/ChangeLog @@ -1,5 +1,11 @@ 2019-05-17 Thomas Schwinge + 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 fc5eb11bb54..be1510a369c 100644 --- 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 index 00000000000..ab0414bfed6 --- /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 } */