Hi! To resolve PR89433 "Repeated use of the OpenACC 'routine' directive" at least for C/C++, I intend to push the attached patches in next GCC development stage 1 (unless that should be addressed right now, and also on other GCC release branches?).
The corresponding Fortran changes will require additional effort. Grüße Thomas
>From 15a0bacff7c75891ceb11a81112b48114c7014a4 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge <[email protected]> Date: Thu, 21 Feb 2019 15:52:20 +0100 Subject: [PATCH 1/3] Refer to OpenACC 'routine' clauses from "omp declare target" attribute TODO gcc/c-family/ * c-attribs.c (c_common_attribute_table): Set min_len to -1 for "omp declare target". gcc/c/ * c-parser.c (c_finish_oacc_routine): Refer to OpenACC 'routine' clauses from "omp declare target" attribute. gcc/cp/ * parser.c (cp_finalize_oacc_routine): Refer to OpenACC 'routine' clauses from "omp declare target" attribute. gcc/testsuite/ * c-c++-common/goacc/classify-routine.c: Update. --- gcc/c-family/c-attribs.c | 2 +- gcc/c/c-parser.c | 2 +- gcc/cp/parser.c | 2 +- gcc/fortran/f95-lang.c | 2 +- gcc/fortran/trans-decl.c | 18 +++++++++++------- .../c-c++-common/goacc/classify-routine.c | 4 ++-- .../gfortran.dg/goacc/classify-routine.f95 | 4 ++-- 7 files changed, 19 insertions(+), 15 deletions(-) diff --git a/gcc/c-family/c-attribs.c b/gcc/c-family/c-attribs.c index 20584593c6a..f108ce6ca6a 100644 --- a/gcc/c-family/c-attribs.c +++ b/gcc/c-family/c-attribs.c @@ -437,7 +437,7 @@ const struct attribute_spec c_common_attribute_table[] = handle_omp_declare_simd_attribute, NULL }, { "simd", 0, 1, true, false, false, false, handle_simd_attribute, NULL }, - { "omp declare target", 0, 0, true, false, false, false, + { "omp declare target", 0, -1, true, false, false, false, handle_omp_declare_target_attribute, NULL }, { "omp declare target link", 0, 0, true, false, false, false, handle_omp_declare_target_attribute, NULL }, diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index 22c7416ac94..c8110faa948 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -15882,7 +15882,7 @@ c_finish_oacc_routine (struct oacc_routine_data *data, tree fndecl, /* Add an "omp declare target" attribute. */ DECL_ATTRIBUTES (fndecl) = tree_cons (get_identifier ("omp declare target"), - NULL_TREE, DECL_ATTRIBUTES (fndecl)); + data->clauses, DECL_ATTRIBUTES (fndecl)); /* Remember that we've used this "#pragma acc routine". */ data->fndecl_seen = true; diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index 545047c91af..418cc507822 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -40181,7 +40181,7 @@ cp_finalize_oacc_routine (cp_parser *parser, tree fndecl, bool is_defn) /* Add an "omp declare target" attribute. */ DECL_ATTRIBUTES (fndecl) = tree_cons (get_identifier ("omp declare target"), - NULL_TREE, DECL_ATTRIBUTES (fndecl)); + parser->oacc_routine->clauses, DECL_ATTRIBUTES (fndecl)); /* Don't unset parser->oacc_routine here: we may still need it to diagnose wrong usage. But, remember that we've used this "#pragma acc diff --git a/gcc/fortran/f95-lang.c b/gcc/fortran/f95-lang.c index ad81f7d93a0..19739792c8f 100644 --- a/gcc/fortran/f95-lang.c +++ b/gcc/fortran/f95-lang.c @@ -91,7 +91,7 @@ static const struct attribute_spec gfc_attribute_table[] = { /* { name, min_len, max_len, decl_req, type_req, fn_type_req, affects_type_identity, handler, exclude } */ - { "omp declare target", 0, 0, true, false, false, false, + { "omp declare target", 0, -1, true, false, false, false, gfc_handle_omp_declare_target_attribute, NULL }, { "omp declare target link", 0, 0, true, false, false, false, gfc_handle_omp_declare_target_attribute, NULL }, diff --git a/gcc/fortran/trans-decl.c b/gcc/fortran/trans-decl.c index 20d453051a2..23b4f98845a 100644 --- a/gcc/fortran/trans-decl.c +++ b/gcc/fortran/trans-decl.c @@ -1400,12 +1400,7 @@ add_attributes_to_decl (symbol_attribute sym_attr, tree list) list = chainon (list, attr); } - if (sym_attr.omp_declare_target_link) - list = tree_cons (get_identifier ("omp declare target link"), - NULL_TREE, list); - else if (sym_attr.omp_declare_target) - list = tree_cons (get_identifier ("omp declare target"), - NULL_TREE, list); + tree clauses = NULL_TREE; if (sym_attr.oacc_routine_lop != OACC_ROUTINE_LOP_NONE) { @@ -1429,11 +1424,20 @@ add_attributes_to_decl (symbol_attribute sym_attr, tree list) gcc_unreachable (); } tree c = build_omp_clause (UNKNOWN_LOCATION, code); + OMP_CLAUSE_CHAIN (c) = clauses; + clauses = c; - tree dims = oacc_build_routine_dims (c); + tree dims = oacc_build_routine_dims (clauses); list = oacc_replace_fn_attrib_attr (list, dims); } + if (sym_attr.omp_declare_target_link) + list = tree_cons (get_identifier ("omp declare target link"), + NULL_TREE, list); + else if (sym_attr.omp_declare_target) + list = tree_cons (get_identifier ("omp declare target"), + clauses, list); + return list; } diff --git a/gcc/testsuite/c-c++-common/goacc/classify-routine.c b/gcc/testsuite/c-c++-common/goacc/classify-routine.c index a723d2cdf51..0b9ba6ea69f 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-routine.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-routine.c @@ -22,10 +22,10 @@ void ROUTINE () } /* Check the offloaded function's attributes. - { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(omp declare target, oacc function \\(0 1, 1 0, 1 0\\)\\)\\)" 1 "ompexp" } } */ + { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(omp declare target \\(worker\\), oacc function \\(0 1, 1 0, 1 0\\)\\)\\)" 1 "ompexp" } } */ /* Check the offloaded function's classification and compute dimensions (will always be 1 x 1 x 1 for non-offloading compilation). { dg-final { scan-tree-dump-times "(?n)Function is OpenACC routine level 1" 1 "oaccdevlow" } } { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccdevlow" } } - { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(0 1, 1 1, 1 1\\), omp declare target, oacc function \\(0 1, 1 0, 1 0\\)\\)\\)" 1 "oaccdevlow" } } */ + { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(0 1, 1 1, 1 1\\), omp declare target \\(worker\\), oacc function \\(0 1, 1 0, 1 0\\)\\)\\)" 1 "oaccdevlow" } } */ diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-routine.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-routine.f95 index e435f5d7eae..401d5270391 100644 --- a/gcc/testsuite/gfortran.dg/goacc/classify-routine.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/classify-routine.f95 @@ -21,10 +21,10 @@ subroutine ROUTINE end subroutine ROUTINE ! Check the offloaded function's attributes. -! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(omp declare target, oacc function \\(0 1, 1 0, 1 0\\)\\)\\)" 1 "ompexp" } } +! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(0 1, 1 0, 1 0\\), omp declare target \\(worker\\)\\)\\)" 1 "ompexp" } } ! Check the offloaded function's classification and compute dimensions (will ! always be 1 x 1 x 1 for non-offloading compilation). ! { dg-final { scan-tree-dump-times "(?n)Function is OpenACC routine level 1" 1 "oaccdevlow" } } ! { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccdevlow" } } -! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(0 1, 1 1, 1 1\\), omp declare target, oacc function \\(0 1, 1 0, 1 0\\)\\)\\)" 1 "oaccdevlow" } } +! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(0 1, 1 1, 1 1\\), omp declare target \\(worker\\)\\)\\)" 1 "oaccdevlow" } } -- 2.17.1
>From a904014e12761dbb861a0a20cbbe03f0b24f2229 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge <[email protected]> Date: Thu, 21 Feb 2019 13:00:34 +0100 Subject: [PATCH 2/3] [C, C++, PR89433] Use 'oacc_verify_routine_clauses' for OpenACC 'routine' directives TODO PR89433 "Repeated use of the OpenACC 'routine' directive" TODO 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-2.c: Update, and move some test into... * c-c++-common/goacc/routine-level-of-parallelism-1.c: ... this new file. --- gcc/c/c-parser.c | 8 + gcc/cp/parser.c | 9 + gcc/omp-general.c | 67 ++++- gcc/omp-general.h | 1 + gcc/testsuite/c-c++-common/goacc/routine-2.c | 20 +- .../goacc/routine-level-of-parallelism-1.c | 265 ++++++++++++++++++ 6 files changed, 341 insertions(+), 29 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 c8110faa948..e2d9e973e09 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -15779,6 +15779,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) { @@ -15793,6 +15796,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 @@ -15856,6 +15862,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 418cc507822..8b02912fec2 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -40025,6 +40025,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 @@ -40121,6 +40124,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. */ @@ -40155,6 +40161,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..07b2df6f8f8 100644 --- a/gcc/omp-general.c +++ b/gcc/omp-general.c @@ -608,9 +608,62 @@ 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; + 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 +685,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<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-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..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
>From 4613129e22af350eb0e6854832582947780e1f64 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge <[email protected]> Date: Thu, 21 Feb 2019 15:42:18 +0100 Subject: [PATCH 3/3] [C, C++, PR89433] Repeated use of the OpenACC 'routine' directive TODO PR89433 - Repeated use of the OpenACC 'routine' directive TODO gcc/ * omp-general.c (oacc_verify_routine_clauses): Change formal parameters. Add checking if already marked as an accelerator routine. Adjust all users. gcc/c/ * c-parser.c (c_finish_oacc_routine): Rework checking if already marked as an accelerator routine. gcc/cp/ * parser.c (cp_finalize_oacc_routine): Rework checking if already marked as an accelerator routine. gcc/testsuite/ * c-c++-common/goacc/oaccdevlow-routine.c: Update. * c-c++-common/goacc/routine-5.c: Likewise. * c-c++-common/goacc/routine-level-of-parallelism-1.c: Likewise. * c-c++-common/goacc/routine-level-of-parallelism-2.c: New file. gcc/testsuite/ * c-c++-common/goacc/routine-1.c: Update. * c-c++-common/goacc/routine-2.c: Likewise. * c-c++-common/goacc/routine-nohost-1.c: Likewise. * g++.dg/goacc/routine-2.C: Likewise. * c-c++-common/goacc/routine-nohost-2.c: New file. --- gcc/c/c-parser.c | 46 ++-- gcc/cp/parser.c | 50 ++-- gcc/omp-general.c | 85 ++++++- gcc/omp-general.h | 3 +- gcc/testsuite/c-c++-common/goacc/routine-5.c | 46 +--- .../goacc/routine-level-of-parallelism-1.c | 236 ++++++++++++++++-- .../goacc/routine-level-of-parallelism-2.c | 71 ++++++ .../goacc/routine-level-of-parallelism-1.f90 | 6 +- 8 files changed, 422 insertions(+), 121 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-2.c diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index e2d9e973e09..04a03fa2feb 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -15862,35 +15862,39 @@ 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)) + int compatible + = oacc_verify_routine_clauses (fndecl, &data->clauses, data->loc, + "#pragma acc routine"); + if (compatible < 0) { - error_at (data->loc, - "%<#pragma acc routine%> already applied to %qD", fndecl); data->error_seen = true; return; } - - if (TREE_USED (fndecl) || (!is_defn && DECL_SAVED_TREE (fndecl))) + if (compatible > 0) { - error_at (data->loc, - TREE_USED (fndecl) - ? G_("%<#pragma acc routine%> must be applied before use") - : G_("%<#pragma acc routine%> must be applied before " - "definition")); - data->error_seen = true; - return; } + else + { + if (TREE_USED (fndecl) || (!is_defn && DECL_SAVED_TREE (fndecl))) + { + error_at (data->loc, + TREE_USED (fndecl) + ? G_("%<#pragma acc routine%> must be applied before use") + : G_("%<#pragma acc routine%> must be applied before" + " definition")); + data->error_seen = true; + return; + } - /* Process the routine's dimension clauses. */ - tree dims = oacc_build_routine_dims (data->clauses); - oacc_replace_fn_attrib (fndecl, dims); + /* Set the routine's level of parallelism. */ + tree dims = oacc_build_routine_dims (data->clauses); + oacc_replace_fn_attrib (fndecl, dims); - /* Add an "omp declare target" attribute. */ - DECL_ATTRIBUTES (fndecl) - = tree_cons (get_identifier ("omp declare target"), - data->clauses, DECL_ATTRIBUTES (fndecl)); + /* Add an "omp declare target" attribute. */ + DECL_ATTRIBUTES (fndecl) + = tree_cons (get_identifier ("omp declare target"), + data->clauses, DECL_ATTRIBUTES (fndecl)); + } /* Remember that we've used this "#pragma acc routine". */ data->fndecl_seen = true; diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index 8b02912fec2..b8a9078ea3b 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -40161,36 +40161,42 @@ 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)) + int compatible + = oacc_verify_routine_clauses (fndecl, &parser->oacc_routine->clauses, + parser->oacc_routine->loc, + "#pragma acc routine"); + if (compatible < 0) { - error_at (parser->oacc_routine->loc, - "%<#pragma acc routine%> already applied to %qD", fndecl); parser->oacc_routine = NULL; return; } - - if (TREE_USED (fndecl) || (!is_defn && DECL_SAVED_TREE (fndecl))) + if (compatible > 0) { - error_at (parser->oacc_routine->loc, - TREE_USED (fndecl) - ? G_("%<#pragma acc routine%> must be applied before use") - : G_("%<#pragma acc routine%> must be applied before " - "definition")); - parser->oacc_routine = NULL; - return; } + else + { + if (TREE_USED (fndecl) || (!is_defn && DECL_SAVED_TREE (fndecl))) + { + error_at (parser->oacc_routine->loc, + TREE_USED (fndecl) + ? G_("%<#pragma acc routine%> must be applied before" + " use") + : G_("%<#pragma acc routine%> must be applied before" + " definition")); + parser->oacc_routine = NULL; + return; + } - /* Process the routine's dimension clauses. */ - tree dims = oacc_build_routine_dims (parser->oacc_routine->clauses); - oacc_replace_fn_attrib (fndecl, dims); + /* Set the routine's level of parallelism. */ + tree dims = oacc_build_routine_dims (parser->oacc_routine->clauses); + oacc_replace_fn_attrib (fndecl, dims); - /* Add an "omp declare target" attribute. */ - DECL_ATTRIBUTES (fndecl) - = tree_cons (get_identifier ("omp declare target"), - parser->oacc_routine->clauses, DECL_ATTRIBUTES (fndecl)); + /* Add an "omp declare target" attribute. */ + DECL_ATTRIBUTES (fndecl) + = tree_cons (get_identifier ("omp declare target"), + parser->oacc_routine->clauses, + DECL_ATTRIBUTES (fndecl)); + } /* Don't unset parser->oacc_routine here: we may still need it to diagnose wrong usage. But, remember that we've used this "#pragma acc diff --git a/gcc/omp-general.c b/gcc/omp-general.c index 07b2df6f8f8..7a7c8ead85c 100644 --- a/gcc/omp-general.c +++ b/gcc/omp-general.c @@ -610,11 +610,14 @@ oacc_set_fn_attrib (tree fn, tree clauses, vec<tree> *args) /* Verify OpenACC routine clauses. - The chain of clauses returned will contain exactly one clause specifying the - level of parallelism. */ + Returns 0 if FNDECL should be marked as an accelerator routine, 1 if it has + already been marked in compatible way, and -1 if incompatible. 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) +int +oacc_verify_routine_clauses (tree fndecl, tree *clauses, location_t loc, + const char *routine_str) { tree c_level = NULL_TREE; tree c_p = NULL_TREE; @@ -660,6 +663,80 @@ oacc_verify_routine_clauses (tree *clauses, location_t loc) OMP_CLAUSE_CHAIN (c_level) = *clauses; *clauses = c_level; } + /* In *clauses, we now have exactly one clause specifying the level of + parallelism. */ + + tree attr + = lookup_attribute ("omp declare target", DECL_ATTRIBUTES (fndecl)); + if (attr != NULL_TREE) + { + /* If a "#pragma acc routine" has already been applied, just verify + this one for compatibility. */ + /* Collect previous directive's clauses. */ + tree c_level_p = NULL_TREE; + for (tree c = TREE_VALUE (attr); 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: + gcc_checking_assert (c_level_p == NULL_TREE); + c_level_p = c; + break; + default: + gcc_unreachable (); + } + gcc_checking_assert (c_level_p != NULL_TREE); + /* ..., and compare to current directive's, which we've already collected + above. */ + tree c_diag; + tree c_diag_p; + /* Matching level of parallelism? */ + if (OMP_CLAUSE_CODE (c_level) != OMP_CLAUSE_CODE (c_level_p)) + { + c_diag = c_level; + c_diag_p = c_level_p; + goto incompatible; + } + /* Compatible. */ + return 1; + + incompatible: + if (c_diag != NULL_TREE) + error_at (OMP_CLAUSE_LOCATION (c_diag), + "incompatible %qs clause when applying" + " %<%s%> to %qD, which has already been" + " marked as an accelerator routine", + omp_clause_code_name[OMP_CLAUSE_CODE (c_diag)], + routine_str, fndecl); + else if (c_diag_p != NULL_TREE) + error_at (loc, + "missing %qs clause when applying" + " %<%s%> to %qD, which has already been" + " marked as an accelerator routine", + omp_clause_code_name[OMP_CLAUSE_CODE (c_diag_p)], + routine_str, fndecl); + else + gcc_unreachable (); + if (c_diag_p != NULL_TREE) + inform (OMP_CLAUSE_LOCATION (c_diag_p), + "... with %qs clause here", + omp_clause_code_name[OMP_CLAUSE_CODE (c_diag_p)]); + else + { + /* In the front ends, we don't preserve location information for the + OpenACC routine directive itself. However, that of c_level_p + should be close. */ + location_t loc_routine = OMP_CLAUSE_LOCATION (c_level_p); + inform (loc_routine, "... without %qs clause near to here", + omp_clause_code_name[OMP_CLAUSE_CODE (c_diag)]); + } + /* Incompatible. */ + return -1; + } + + return 0; } /* Process the OpenACC routine's clauses to generate an attribute diff --git a/gcc/omp-general.h b/gcc/omp-general.h index 4241c33d99e..f96d3c7768a 100644 --- a/gcc/omp-general.h +++ b/gcc/omp-general.h @@ -84,7 +84,8 @@ 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 int oacc_verify_routine_clauses (tree, tree *, location_t, + const char *); 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-5.c b/gcc/testsuite/c-c++-common/goacc/routine-5.c index b967a7447bd..a8c974cb6bf 100644 --- a/gcc/testsuite/c-c++-common/goacc/routine-5.c +++ b/gcc/testsuite/c-c++-common/goacc/routine-5.c @@ -150,61 +150,19 @@ void f_static_assert(); #pragma acc routine __extension__ extern void ex1(); -#pragma acc routine (ex1) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*ex1" } */ +#pragma acc routine (ex1) worker /* { dg-error "has already been marked as an accelerator routine" } */ #pragma acc routine __extension__ __extension__ __extension__ __extension__ __extension__ void ex2() { } -#pragma acc routine (ex2) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*ex2" } */ +#pragma acc routine (ex2) worker /* { dg-error "has already been marked as an accelerator routine" } */ #pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ __extension__ int ex3; #pragma acc routine (ex3) /* { dg-error ".ex3. does not refer to a function" } */ -/* "#pragma acc routine" already applied. */ - -extern void fungsi_1(); -#pragma acc routine(fungsi_1) gang -#pragma acc routine(fungsi_1) gang /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_1" } */ -#pragma acc routine(fungsi_1) worker /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_1" } */ -#pragma acc routine(fungsi_1) vector /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_1" } */ - -#pragma acc routine seq -extern void fungsi_2(); -#pragma acc routine(fungsi_2) seq /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_2." } */ -#pragma acc routine(fungsi_2) worker /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_2." } */ -#pragma acc routine(fungsi_2) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_2." } */ - -#pragma acc routine vector -extern void fungsi_3(); -#pragma acc routine vector /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_3." } */ -void fungsi_3() -{ -} - -extern void fungsi_4(); -#pragma acc routine (fungsi_4) worker -#pragma acc routine gang /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_4." } */ -void fungsi_4() -{ -} - -#pragma acc routine gang -void fungsi_5() -{ -} -#pragma acc routine (fungsi_5) worker /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_5." } */ - -#pragma acc routine seq -void fungsi_6() -{ -} -#pragma acc routine seq /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_6." } */ -extern void fungsi_6(); - - /* "#pragma acc routine" must be applied before. */ void Bar (); 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 index d71d6e088f7..aa3c1243ba1 100644 --- 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 @@ -1,6 +1,5 @@ /* 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. */ + parallelism with the OpenACC routine directive. */ extern void g_1 (void); #pragma acc routine (g_1) gang gang /* { dg-error "too many 'gang' clauses" } */ @@ -42,10 +41,10 @@ void s_2 (void) void g_3 (void) { } -#pragma acc routine (g_3) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*g_3." } */ \ +#pragma acc routine (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." } */ \ +#pragma acc routine (g_3) \ gang \ vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ @@ -53,10 +52,10 @@ 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." } */ \ +#pragma acc routine (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." } */ \ +#pragma acc routine (w_3) \ worker \ seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */ @@ -66,10 +65,10 @@ extern void w_3 (void); void v_3 (void) { } -#pragma acc routine (v_3) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*v_3." } */ \ +#pragma acc routine (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." } */ \ +#pragma acc routine (v_3) \ vector \ gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */ @@ -77,10 +76,10 @@ 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." } */ \ +#pragma acc routine (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." } */ \ +#pragma acc routine (s_3) \ seq \ worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */ @@ -91,12 +90,12 @@ extern void s_3 (void); 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." } */ \ +#pragma acc routine (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." } */ \ +#pragma acc routine (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" } */ \ @@ -108,12 +107,12 @@ extern void w_4 (void); 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." } */ \ +#pragma acc routine (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." } */ \ +#pragma acc routine (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" } */ \ @@ -127,12 +126,12 @@ extern void w_4 (void); void v_4 (void) { } -#pragma acc routine (v_4) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*v_4." } */ \ +#pragma acc routine (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." } */ \ +#pragma acc routine (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" } */ \ @@ -146,12 +145,12 @@ void v_4 (void) void s_4 (void) { } -#pragma acc routine (s_4) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*s_4." } */ \ +#pragma acc routine (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." } */ \ +#pragma acc routine (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" } */ \ @@ -169,7 +168,7 @@ void s_4 (void) void g_5 (void) { } -#pragma acc routine (g_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*g_5." } */ \ +#pragma acc routine (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 } */ \ @@ -177,7 +176,7 @@ void g_5 (void) /* { 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." } */ \ +#pragma acc routine (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 } */ \ @@ -195,7 +194,7 @@ void g_5 (void) 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." } */ \ +#pragma acc routine (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 } */ \ @@ -203,7 +202,7 @@ extern void w_5 (void); /* { 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." } */ \ +#pragma acc routine (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 } */ \ @@ -221,7 +220,7 @@ extern void w_5 (void); 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." } */ \ +#pragma acc routine (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 } */ \ @@ -229,7 +228,7 @@ extern void v_5 (void); /* { 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." } */ \ +#pragma acc routine (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 } */ \ @@ -247,7 +246,7 @@ extern void s_5 (void); /* { 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." } */ \ +#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 } */ \ @@ -255,7 +254,7 @@ extern void s_5 (void); /* { 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." } */ \ +#pragma acc routine (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 } */ \ @@ -263,3 +262,188 @@ extern void s_5 (void); /* { 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 } */ + + +/* Like the *_5 tests, but with the order of clauses changed in the second and + following routine directives for the specific *_5 function. */ + +#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_6 (void) +{ +} +#pragma acc routine (g_6) \ + vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \ + /* { dg-error "incompatible .vector. clause when applying .#pragma acc routine. to .\[void \]*g_6\[\\(\\)\]*., which has already been marked as an accelerator routine" "" { target *-*-* } .-1 } */ \ + gang 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 } */ \ + 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_6) \ + seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \ + /* { dg-error "incompatible .seq. clause when applying .#pragma acc routine. to .\[void \]*g_6\[\\(\\)\]*., which has already been marked as an accelerator routine" "" { target *-*-* } .-1 } */ \ + gang 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 } */ \ + 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_6 (void); +#pragma acc routine (w_6) \ + vector vector vector /* { dg-error "too many 'vector' clauses" } */ \ + /* { dg-error "incompatible .vector. clause when applying .#pragma acc routine. to .\[void \]*w_6\[\\(\\)\]*., which has already been marked as an accelerator routine" "" { 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 /* { 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_6) \ + seq seq /* { dg-error "too many 'seq' clauses" } */ \ + /* { dg-error "incompatible .seq. clause when applying .#pragma acc routine. to .\[void \]*w_6\[\\(\\)\]*., which has already been marked as an accelerator routine" "" { target *-*-* } .-1 } */ \ + 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 } */ \ + 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_6 (void); +#pragma acc routine (v_6) \ + seq seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \ + /* { dg-error "incompatible .seq. clause when applying .#pragma acc routine. to .\[void \]*v_6\[\\(\\)\]*., which has already been marked as an accelerator routine" "" { target *-*-* } .-1 } */ \ + vector vector vector /* { dg-error "too many 'vector' clauses" } */ \ + /* { dg-error ".vector. 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_6) \ + gang gang /* { dg-error "too many 'gang' clauses" } */ \ + /* { dg-error "incompatible .gang. clause when applying .#pragma acc routine. to .\[void \]*v_6\[\\(\\)\]*., which has already been marked as an accelerator routine" "" { 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 } */ \ + worker worker /* { dg-error "too many 'worker' clauses" } */ \ + /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ + +extern void s_6 (void); +#pragma acc routine (s_6) \ + 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_6) \ + vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \ + /* { dg-error "incompatible .vector. clause when applying .#pragma acc routine. to .\[void \]*s_6\[\\(\\)\]*., which has already been marked as an accelerator routine" "" { 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 } */ \ + worker worker worker /* { dg-error "too many 'worker' clauses" } */ \ + /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ +#pragma acc routine (s_6) \ + worker worker /* { dg-error "too many 'worker' clauses" } */ \ + /* { dg-error "incompatible .worker. clause when applying .#pragma acc routine. to .\[void \]*s_6\[\\(\\)\]*., which has already been marked as an accelerator routine" "" { target *-*-* } .-1 } */ \ + 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 } */ \ + gang gang /* { dg-error "too many 'gang' clauses" } */ \ + /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ + + +/* Like the *_6 tests, but without all the duplicate clauses, so that the + routine directives are valid in isolation. */ + +#pragma acc routine \ + gang +void g_7 (void) +{ +} +#pragma acc routine (g_7) \ + vector /* { dg-error "incompatible .vector. clause when applying .#pragma acc routine. to .\[void \]*g_7\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ +#pragma acc routine (g_7) \ + seq /* { dg-error "incompatible .seq. clause when applying .#pragma acc routine. to .\[void \]*g_7\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ + +#pragma acc routine \ + worker +extern void w_7 (void); +#pragma acc routine (w_7) \ + vector /* { dg-error "incompatible .vector. clause when applying .#pragma acc routine. to .\[void \]*w_7\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ +#pragma acc routine (w_7) \ + seq /* { dg-error "incompatible .seq. clause when applying .#pragma acc routine. to .\[void \]*w_7\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ + +#pragma acc routine \ + vector +extern void v_7 (void); +#pragma acc routine (v_7) \ + seq /* { dg-error "incompatible .seq. clause when applying .#pragma acc routine. to .\[void \]*v_7\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ +#pragma acc routine (v_7) \ + gang /* { dg-error "incompatible .gang. clause when applying .#pragma acc routine. to .\[void \]*v_7\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ + +extern void s_7 (void); +#pragma acc routine (s_7) \ + seq +#pragma acc routine (s_7) \ + vector /* { dg-error "incompatible .vector. clause when applying .#pragma acc routine. to .\[void \]*s_7\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ +#pragma acc routine (s_7) \ + worker /* { dg-error "incompatible .worker. clause when applying .#pragma acc routine. to .\[void \]*s_7\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ + + +/* Test cases for implicit seq clause. */ + +#pragma acc routine \ + gang +void g_8 (void) +{ +} +#pragma acc routine (g_8) /* { dg-error "incompatible .seq. clause when applying .#pragma acc routine. to .\[void \]*g_8\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ + +#pragma acc routine \ + worker +extern void w_8 (void); +#pragma acc routine (w_8) /* { dg-error "incompatible .seq. clause when applying .#pragma acc routine. to .\[void \]*w_8\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ + +#pragma acc routine \ + vector +extern void v_8 (void); +#pragma acc routine (v_8) /* { dg-error "incompatible .seq. clause when applying .#pragma acc routine. to .\[void \]*v_8\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ + +extern void s_8 (void); +#pragma acc routine (s_8) +#pragma acc routine (s_8) \ + vector /* { dg-error "incompatible .vector. clause when applying .#pragma acc routine. to .\[void \]*s_8\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ +#pragma acc routine (s_8) \ + gang /* { dg-error "incompatible .gang. clause when applying .#pragma acc routine. to .\[void \]*s_8\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ +#pragma acc routine (s_8) \ + worker /* { dg-error "incompatible .worker. clause when applying .#pragma acc routine. to .\[void \]*s_8\[\\(\\)\]*., which has already been marked as an accelerator routine" } */ diff --git a/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-2.c b/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-2.c new file mode 100644 index 00000000000..db4360a81ec --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-2.c @@ -0,0 +1,71 @@ +/* Test various aspects of clauses specifying compatible levels of parallelism + with the OpenACC routine directive. The Fortran counterpart is + '../../gfortran.dg/goacc/routine-level-of-parallelism-1.f90'. */ + +#pragma acc routine gang +void g_1 (void) +{ +} +#pragma acc routine (g_1) gang +#pragma acc routine (g_1) gang + + +extern void w_1 (void); +#pragma acc routine (w_1) worker +#pragma acc routine (w_1) worker +#pragma acc routine (w_1) worker + + +#pragma acc routine vector +extern void v_1 (void); +#pragma acc routine (v_1) vector +#pragma acc routine (v_1) vector + + +/* Also test the implicit seq clause. */ + +#pragma acc routine seq +extern void s_1_1 (void); +#pragma acc routine (s_1_1) +#pragma acc routine (s_1_1) seq +#pragma acc routine (s_1_1) +#pragma acc routine (s_1_1) seq + +#pragma acc routine +extern void s_1_2 (void); +#pragma acc routine (s_1_2) +#pragma acc routine (s_1_2) seq +#pragma acc routine (s_1_2) +#pragma acc routine (s_1_2) seq + +extern void s_2_1 (void); +#pragma acc routine (s_2_1) seq +#pragma acc routine (s_2_1) +#pragma acc routine (s_2_1) seq +#pragma acc routine (s_2_1) +#pragma acc routine (s_2_1) seq + +extern void s_2_2 (void); +#pragma acc routine (s_2_2) +#pragma acc routine (s_2_2) +#pragma acc routine (s_2_2) seq +#pragma acc routine (s_2_2) +#pragma acc routine (s_2_2) seq + +#pragma acc routine seq +void s_3_1 (void) +{ +} +#pragma acc routine (s_3_1) +#pragma acc routine (s_3_1) seq +#pragma acc routine (s_3_1) +#pragma acc routine (s_3_1) seq + +#pragma acc routine +void s_3_2 (void) +{ +} +#pragma acc routine (s_3_2) +#pragma acc routine (s_3_2) seq +#pragma acc routine (s_3_2) +#pragma acc routine (s_3_2) seq diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-level-of-parallelism-1.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-level-of-parallelism-1.f90 index 75dd1b01f6f..83b8c24b41d 100644 --- a/gcc/testsuite/gfortran.dg/goacc/routine-level-of-parallelism-1.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/routine-level-of-parallelism-1.f90 @@ -1,6 +1,6 @@ -! Test various aspects of clauses specifying compatible levels of -! parallelism with the OpenACC routine directive. The Fortran counterpart is -! c-c++-common/goacc/routine-level-of-parallelism-2.c +! Test various aspects of clauses specifying compatible levels of parallelism +! with the OpenACC routine directive. The C/C++ counterpart is +! '../../c-c++-common/goacc/routine-level-of-parallelism-2.c'. subroutine g_1 !$acc routine gang -- 2.17.1
