From 3608ceecdf0930a55cfa65a4e8ba534278bca0e0 Mon Sep 17 00:00:00 2001 From: Julian Brown Date: Thu, 7 Feb 2019 05:19:53 -0800 Subject: [PATCH] Use oacc_verify_routine_clauses for C/C++ 2018-10-02 Thomas Schwinge Cesar Philippidis 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) --- gcc/ChangeLog.omp | 8 + gcc/c/ChangeLog.omp | 6 + gcc/c/c-parser.c | 8 + gcc/cp/ChangeLog.omp | 7 + gcc/cp/parser.c | 9 + gcc/omp-general.c | 69 ++++- gcc/omp-general.h | 1 + gcc/testsuite/ChangeLog.omp | 5 + .../goacc/routine-level-of-parallelism-1.c | 265 ++++++++++++++++++ 9 files changed, 368 insertions(+), 10 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-1.c diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp index d1e1ab54ea9e..d50b791fa963 100644 --- a/gcc/ChangeLog.omp +++ b/gcc/ChangeLog.omp @@ -1,3 +1,11 @@ +2018-10-02 Thomas Schwinge + Cesar Philippidis + + * 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 Cesar Philippidis diff --git a/gcc/c/ChangeLog.omp b/gcc/c/ChangeLog.omp index 023607f1e4d9..ed7b8b9e73e7 100644 --- a/gcc/c/ChangeLog.omp +++ b/gcc/c/ChangeLog.omp @@ -1,3 +1,9 @@ +2018-10-02 Thomas Schwinge + Cesar Philippidis + + * 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 Cesar Philippidis diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index 60e4ed5b63b3..7c2907a765c2 100644 --- 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 f9ba662d4aa6..63bae9714ff0 100644 --- a/gcc/cp/ChangeLog.omp +++ b/gcc/cp/ChangeLog.omp @@ -1,3 +1,10 @@ +2018-10-02 Thomas Schwinge + Cesar Philippidis + + * 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 Cesar Philippidis diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index 418be26be4b6..4e3993c0f69f 100644 --- 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 356772ff4582..b3cd9cb103d9 100644 --- a/gcc/omp-general.c +++ b/gcc/omp-general.c @@ -608,9 +608,64 @@ 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. + + 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 60faa5213a27..4241c33d99e0 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.omp b/gcc/testsuite/ChangeLog.omp index 2a0d2f461580..adad65edb20f 100644 --- a/gcc/testsuite/ChangeLog.omp +++ b/gcc/testsuite/ChangeLog.omp @@ -1,3 +1,8 @@ +2018-10-02 Thomas Schwinge + Cesar Philippidis + + * c-c++-common/goacc/routine-level-of-parallelism-1.c: New test. + 2018-10-02 Thomas Schwinge Cesar Philippidis 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 000000000000..d71d6e088f79 --- /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.47.2