Hi!
On Fri, 22 Feb 2019 12:22:38 +0100, I wrote:
> 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?).
(I still ponder on the question whether to eventually backport all this
to other GCC release branches...)
> The corresponding Fortran changes will require additional effort.
..., which I actually already have committed a while ago.
Now committed to trunk in r271343 "[PR89433] Refer to OpenACC 'routine'
clauses from "omp declare target" attribute", r271344 "[PR89433] Use
'oacc_verify_routine_clauses' for C/C++ OpenACC 'routine' directives",
r271345 "[PR89433] Repeated use of the C/C++ OpenACC 'routine'
directive", see attached.
Grüße
Thomas
>From 9b5009857b50944d7711f34dd0ac98364b7f1f9b Mon Sep 17 00:00:00 2001
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Fri, 17 May 2019 19:13:04 +0000
Subject: [PATCH 1/3] [PR89433] Refer to OpenACC 'routine' clauses from "omp
declare target" attribute
gcc/c-family/
PR c/89433
* c-attribs.c (c_common_attribute_table): Set min_len to -1 for
"omp declare target".
gcc/c/
PR c/89433
* c-parser.c (c_finish_oacc_routine): Refer to OpenACC 'routine'
clauses from "omp declare target" attribute.
gcc/cp/
PR c++/89433
* parser.c (cp_finalize_oacc_routine): Refer to OpenACC 'routine'
clauses from "omp declare target" attribute.
gcc/fortran/
PR fortran/89433
* f95-lang.c (gfc_attribute_table): Set min_len to -1 for "omp
declare target".
* trans-decl.c (add_attributes_to_decl): Refer to OpenACC
'routine' clauses from "omp declare target" attribute.
gcc/testsuite/
PR testsuite/89433
* c-c++-common/goacc/classify-routine.c: Update.
* gfortran.dg/goacc/classify-routine.f95: Likewise.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@271343 138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/c-family/ChangeLog | 6 ++++++
gcc/c-family/c-attribs.c | 2 +-
gcc/c/ChangeLog | 6 ++++++
gcc/c/c-parser.c | 2 +-
gcc/cp/ChangeLog | 6 ++++++
gcc/cp/parser.c | 2 +-
gcc/fortran/ChangeLog | 8 ++++++++
gcc/fortran/f95-lang.c | 2 +-
gcc/fortran/trans-decl.c | 18 +++++++++++-------
gcc/testsuite/ChangeLog | 6 ++++++
.../c-c++-common/goacc/classify-routine.c | 4 ++--
.../gfortran.dg/goacc/classify-routine.f95 | 4 ++--
12 files changed, 51 insertions(+), 15 deletions(-)
diff --git a/gcc/c-family/ChangeLog b/gcc/c-family/ChangeLog
index 47c1d3d51f50..50687764221b 100644
--- a/gcc/c-family/ChangeLog
+++ b/gcc/c-family/ChangeLog
@@ -1,3 +1,9 @@
+2019-05-17 Thomas Schwinge <tho...@codesourcery.com>
+
+ PR c/89433
+ * c-attribs.c (c_common_attribute_table): Set min_len to -1 for
+ "omp declare target".
+
2019-05-16 Martin Sebor <mse...@redhat.com>
* c-attribs.c (handle_no_sanitize_attribute): Quote identifiers,
diff --git a/gcc/c-family/c-attribs.c b/gcc/c-family/c-attribs.c
index 12c0b9bfb543..03203470955a 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/ChangeLog b/gcc/c/ChangeLog
index f558d28dc259..f0cab2e65929 100644
--- a/gcc/c/ChangeLog
+++ b/gcc/c/ChangeLog
@@ -1,3 +1,9 @@
+2019-05-17 Thomas Schwinge <tho...@codesourcery.com>
+
+ PR c/89433
+ * c-parser.c (c_finish_oacc_routine): Refer to OpenACC 'routine'
+ clauses from "omp declare target" attribute.
+
2019-05-16 Martin Sebor <mse...@redhat.com>
* c-decl.c (start_decl): Quote keywords, operators, and
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 993cfe05ecb0..3cbbb199bdda 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -15904,7 +15904,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/ChangeLog b/gcc/cp/ChangeLog
index 08b7d5337071..40622acc0ff2 100644
--- a/gcc/cp/ChangeLog
+++ b/gcc/cp/ChangeLog
@@ -1,3 +1,9 @@
+2019-05-17 Thomas Schwinge <tho...@codesourcery.com>
+
+ PR c++/89433
+ * parser.c (cp_finalize_oacc_routine): Refer to OpenACC 'routine'
+ clauses from "omp declare target" attribute.
+
2019-05-16 Martin Sebor <mse...@redhat.com>
* call.c (print_z_candidate): Wrap diagnostic text in a gettext
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index e6ef5a9bc008..15424b6b6337 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -40292,7 +40292,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/ChangeLog b/gcc/fortran/ChangeLog
index b8851067cf3e..f09e715353bd 100644
--- a/gcc/fortran/ChangeLog
+++ b/gcc/fortran/ChangeLog
@@ -1,3 +1,11 @@
+2019-05-17 Thomas Schwinge <tho...@codesourcery.com>
+
+ PR fortran/89433
+ * f95-lang.c (gfc_attribute_table): Set min_len to -1 for "omp
+ declare target".
+ * trans-decl.c (add_attributes_to_decl): Refer to OpenACC
+ 'routine' clauses from "omp declare target" attribute.
+
2019-05-16 Martin Sebor <mse...@redhat.com>
* gfortranspec.c (append_arg): Spell out the word "argument."
diff --git a/gcc/fortran/f95-lang.c b/gcc/fortran/f95-lang.c
index 3e3d3046bdb5..6b9f490d2bbc 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 07d1c33af72c..8420870a6b78 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)
{
@@ -1430,11 +1425,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/ChangeLog b/gcc/testsuite/ChangeLog
index e1aa25736efb..0f44c6d4ab0e 100644
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,3 +1,9 @@
+2019-05-17 Thomas Schwinge <tho...@codesourcery.com>
+
+ PR testsuite/89433
+ * c-c++-common/goacc/classify-routine.c: Update.
+ * gfortran.dg/goacc/classify-routine.f95: Likewise.
+
2019-05-16 Martin Sebor <mse...@redhat.com>
* c-c++-common/Wbool-operation-1.c: Adjust text of expected diagnostics.
diff --git a/gcc/testsuite/c-c++-common/goacc/classify-routine.c b/gcc/testsuite/c-c++-common/goacc/classify-routine.c
index a723d2cdf513..0b9ba6ea69fc 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 e435f5d7eaef..401d5270391e 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 5f7ea2eebbef8715d22695470221dff57fe45c42 Mon Sep 17 00:00:00 2001
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Fri, 17 May 2019 19:13:15 +0000
Subject: [PATCH 2/3] [PR89433] Use 'oacc_verify_routine_clauses' for C/C++
OpenACC 'routine' directives
gcc/
PR middle-end/89433
* omp-general.c (oacc_build_routine_dims): Move some of its
processing into...
(oacc_verify_routine_clauses): ... this new function.
* omp-general.h (oacc_verify_routine_clauses): New prototype.
gcc/c/
PR c/89433
* c-parser.c (c_parser_oacc_routine): Normalize order of clauses.
(c_finish_oacc_routine): Call oacc_verify_routine_clauses.
gcc/cp/
PR c++/89433
* parser.c (cp_parser_oacc_routine)
(cp_parser_late_parsing_oacc_routine): Normalize order of clauses.
(cp_finalize_oacc_routine): Call oacc_verify_routine_clauses.
gcc/testsuite/
PR testsuite/89433
* c-c++-common/goacc/routine-2.c: Update, and move some test
into...
* c-c++-common/goacc/routine-level-of-parallelism-1.c: ... this
new file.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@271344 138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/ChangeLog | 8 +
gcc/c/ChangeLog | 4 +
gcc/c/c-parser.c | 8 +
gcc/cp/ChangeLog | 5 +
gcc/cp/parser.c | 9 +
gcc/omp-general.c | 66 ++++-
gcc/omp-general.h | 1 +
gcc/testsuite/ChangeLog | 6 +
gcc/testsuite/c-c++-common/goacc/routine-2.c | 20 +-
.../goacc/routine-level-of-parallelism-1.c | 264 ++++++++++++++++++
10 files changed, 362 insertions(+), 29 deletions(-)
create mode 100644 gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-1.c
diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index 5fc97248662d..b8a435860525 100644
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,11 @@
+2019-05-17 Thomas Schwinge <tho...@codesourcery.com>
+
+ PR middle-end/89433
+ * omp-general.c (oacc_build_routine_dims): Move some of its
+ processing into...
+ (oacc_verify_routine_clauses): ... this new function.
+ * omp-general.h (oacc_verify_routine_clauses): New prototype.
+
2019-05-17 Iain Sandoe <i...@sandoe.co.uk>
* config/rs6000/rs6000.c (machopic_output_stub): Adjust the
diff --git a/gcc/c/ChangeLog b/gcc/c/ChangeLog
index f0cab2e65929..1393e8f47fd2 100644
--- a/gcc/c/ChangeLog
+++ b/gcc/c/ChangeLog
@@ -1,5 +1,9 @@
2019-05-17 Thomas Schwinge <tho...@codesourcery.com>
+ PR c/89433
+ * c-parser.c (c_parser_oacc_routine): Normalize order of clauses.
+ (c_finish_oacc_routine): Call oacc_verify_routine_clauses.
+
PR c/89433
* c-parser.c (c_finish_oacc_routine): Refer to OpenACC 'routine'
clauses from "omp declare target" attribute.
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 3cbbb199bdda..8337f1cce0cf 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -15801,6 +15801,9 @@ c_parser_oacc_routine (c_parser *parser, enum pragma_context context)
data.clauses
= c_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK,
"#pragma acc routine");
+ /* The clauses are in reverse order; fix that to make later diagnostic
+ emission easier. */
+ data.clauses = nreverse (data.clauses);
if (TREE_CODE (decl) != FUNCTION_DECL)
{
@@ -15815,6 +15818,9 @@ c_parser_oacc_routine (c_parser *parser, enum pragma_context context)
data.clauses
= c_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK,
"#pragma acc routine");
+ /* The clauses are in reverse order; fix that to make later diagnostic
+ emission easier. */
+ data.clauses = nreverse (data.clauses);
/* Emit a helpful diagnostic if there's another pragma following this
one. Also don't allow a static assertion declaration, as in the
@@ -15878,6 +15884,8 @@ c_finish_oacc_routine (struct oacc_routine_data *data, tree fndecl,
return;
}
+ oacc_verify_routine_clauses (&data->clauses, data->loc);
+
if (oacc_get_fn_attrib (fndecl))
{
error_at (data->loc,
diff --git a/gcc/cp/ChangeLog b/gcc/cp/ChangeLog
index 40622acc0ff2..2f1e06ca458f 100644
--- a/gcc/cp/ChangeLog
+++ b/gcc/cp/ChangeLog
@@ -1,5 +1,10 @@
2019-05-17 Thomas Schwinge <tho...@codesourcery.com>
+ PR c++/89433
+ * parser.c (cp_parser_oacc_routine)
+ (cp_parser_late_parsing_oacc_routine): Normalize order of clauses.
+ (cp_finalize_oacc_routine): Call oacc_verify_routine_clauses.
+
PR c++/89433
* parser.c (cp_finalize_oacc_routine): Refer to OpenACC 'routine'
clauses from "omp declare target" attribute.
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 15424b6b6337..aa6507e42f41 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -40136,6 +40136,9 @@ cp_parser_oacc_routine (cp_parser *parser, cp_token *pragma_tok,
= cp_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK,
"#pragma acc routine",
cp_lexer_peek_token (parser->lexer));
+ /* The clauses are in reverse order; fix that to make later diagnostic
+ emission easier. */
+ data.clauses = nreverse (data.clauses);
if (decl && is_overloaded_fn (decl)
&& (TREE_CODE (decl) != FUNCTION_DECL
@@ -40232,6 +40235,9 @@ cp_parser_late_parsing_oacc_routine (cp_parser *parser, tree attrs)
parser->oacc_routine->clauses
= cp_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK,
"#pragma acc routine", pragma_tok);
+ /* The clauses are in reverse order; fix that to make later diagnostic
+ emission easier. */
+ parser->oacc_routine->clauses = nreverse (parser->oacc_routine->clauses);
cp_parser_pop_lexer (parser);
/* Later, cp_finalize_oacc_routine will process the clauses, and then set
fndecl_seen. */
@@ -40266,6 +40272,9 @@ cp_finalize_oacc_routine (cp_parser *parser, tree fndecl, bool is_defn)
return;
}
+ oacc_verify_routine_clauses (&parser->oacc_routine->clauses,
+ parser->oacc_routine->loc);
+
if (oacc_get_fn_attrib (fndecl))
{
error_at (parser->oacc_routine->loc,
diff --git a/gcc/omp-general.c b/gcc/omp-general.c
index 356772ff4582..f1d859b0275c 100644
--- a/gcc/omp-general.c
+++ b/gcc/omp-general.c
@@ -608,9 +608,61 @@ 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.
+
+ Upon returning, the chain of clauses will contain exactly one clause
+ specifying the level of parallelism. */
+
+void
+oacc_verify_routine_clauses (tree *clauses, location_t loc)
+{
+ tree c_level = NULL_TREE;
+ tree c_p = NULL_TREE;
+ for (tree c = *clauses; c; c_p = c, c = OMP_CLAUSE_CHAIN (c))
+ switch (OMP_CLAUSE_CODE (c))
+ {
+ case OMP_CLAUSE_GANG:
+ case OMP_CLAUSE_WORKER:
+ case OMP_CLAUSE_VECTOR:
+ case OMP_CLAUSE_SEQ:
+ if (c_level == NULL_TREE)
+ c_level = c;
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_CODE (c_level))
+ {
+ /* This has already been diagnosed in the front ends. */
+ /* Drop the duplicate clause. */
+ gcc_checking_assert (c_p != NULL_TREE);
+ OMP_CLAUSE_CHAIN (c_p) = OMP_CLAUSE_CHAIN (c);
+ c = c_p;
+ }
+ else
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%qs specifies a conflicting level of parallelism",
+ omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+ inform (OMP_CLAUSE_LOCATION (c_level),
+ "... to the previous %qs clause here",
+ omp_clause_code_name[OMP_CLAUSE_CODE (c_level)]);
+ /* Drop the conflicting clause. */
+ gcc_checking_assert (c_p != NULL_TREE);
+ OMP_CLAUSE_CHAIN (c_p) = OMP_CLAUSE_CHAIN (c);
+ c = c_p;
+ }
+ break;
+ default:
+ gcc_unreachable ();
+ }
+ if (c_level == NULL_TREE)
+ {
+ /* Default to an implicit 'seq' clause. */
+ c_level = build_omp_clause (loc, OMP_CLAUSE_SEQ);
+ OMP_CLAUSE_CHAIN (c_level) = *clauses;
+ *clauses = c_level;
+ }
+}
+
+/* Process the OpenACC 'routine' directive clauses to generate an attribute
+ for the level of parallelism. All dimensions have a size of zero
(dynamic). TREE_PURPOSE is set to indicate whether that dimension
can have a loop partitioned on it. non-zero indicates
yes, zero indicates no. By construction once a non-zero has been
@@ -632,16 +684,10 @@ oacc_build_routine_dims (tree clauses)
for (ix = GOMP_DIM_MAX + 1; ix--;)
if (OMP_CLAUSE_CODE (clauses) == ids[ix])
{
- if (level >= 0)
- error_at (OMP_CLAUSE_LOCATION (clauses),
- "multiple loop axes specified for routine");
level = ix;
break;
}
-
- /* Default to SEQ. */
- if (level < 0)
- level = GOMP_DIM_MAX;
+ gcc_checking_assert (level >= 0);
tree dims = NULL_TREE;
diff --git a/gcc/omp-general.h b/gcc/omp-general.h
index 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<tree> *args);
+extern void oacc_verify_routine_clauses (tree *, location_t);
extern tree oacc_build_routine_dims (tree clauses);
extern tree oacc_get_fn_attrib (tree fn);
extern bool offloading_function_p (tree fn);
diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog
index 0f44c6d4ab0e..4b07888dadbc 100644
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,5 +1,11 @@
2019-05-17 Thomas Schwinge <tho...@codesourcery.com>
+ PR testsuite/89433
+ * c-c++-common/goacc/routine-2.c: Update, and move some test
+ into...
+ * c-c++-common/goacc/routine-level-of-parallelism-1.c: ... this
+ new file.
+
PR testsuite/89433
* c-c++-common/goacc/classify-routine.c: Update.
* gfortran.dg/goacc/classify-routine.f95: Likewise.
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-2.c b/gcc/testsuite/c-c++-common/goacc/routine-2.c
index fc5eb11bb54d..be1510a369ca 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 000000000000..ab0414bfed69
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-1.c
@@ -0,0 +1,264 @@
+/* Test various aspects of clauses specifying incompatible levels of
+ parallelism with the OpenACC 'routine' directive. */
+
+extern void g_1 (void);
+#pragma acc routine (g_1) gang gang /* { dg-error "too many 'gang' clauses" } */
+
+#pragma acc routine worker worker /* { dg-error "too many 'worker' clauses" } */
+void w_1 (void)
+{
+}
+
+#pragma acc routine vector vector /* { dg-error "too many 'vector' clauses" } */
+void v_1 (void)
+{
+}
+
+#pragma acc routine seq seq /* { dg-error "too many 'seq' clauses" } */
+extern void s_1 (void);
+
+
+#pragma acc routine gang gang gang /* { dg-error "too many 'gang' clauses" } */
+void g_2 (void)
+{
+}
+
+#pragma acc routine worker worker worker /* { dg-error "too many 'worker' clauses" } */
+extern void w_2 (void);
+
+extern void v_2 (void);
+#pragma acc routine (v_2) vector vector vector /* { dg-error "too many 'vector' clauses" } */
+
+#pragma acc routine seq seq seq /* { dg-error "too many 'seq' clauses" } */
+void s_2 (void)
+{
+}
+
+
+#pragma acc routine \
+ gang \
+ worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */
+void g_3 (void)
+{
+}
+#pragma acc routine (g_3) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*g_3." } */ \
+ gang \
+ seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */
+#pragma acc routine (g_3) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*g_3." } */ \
+ gang \
+ vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */
+
+extern void w_3 (void);
+#pragma acc routine (w_3) \
+ worker \
+ vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */
+#pragma acc routine (w_3) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*w_3." } */ \
+ worker \
+ gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */
+#pragma acc routine (w_3) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*w_3." } */ \
+ worker \
+ seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */
+
+#pragma acc routine \
+ vector \
+ seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */
+void v_3 (void)
+{
+}
+#pragma acc routine (v_3) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*v_3." } */ \
+ vector \
+ worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */
+#pragma acc routine (v_3) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*v_3." } */ \
+ vector \
+ gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */
+
+extern void s_3 (void);
+#pragma acc routine (s_3) \
+ seq \
+ gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */
+#pragma acc routine (s_3) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*s_3." } */ \
+ seq \
+ vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */
+#pragma acc routine (s_3) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*s_3." } */ \
+ seq \
+ worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */
+
+
+#pragma acc routine \
+ gang gang gang /* { dg-error "too many 'gang' clauses" } */ \
+ worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */ \
+ vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ \
+ seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */
+extern void g_4 (void);
+#pragma acc routine (g_4) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*g_4." } */ \
+ gang gang gang /* { dg-error "too many 'gang' clauses" } */ \
+ vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ \
+ worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */ \
+ seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */
+#pragma acc routine (g_4) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*g_4." } */ \
+ gang gang gang /* { dg-error "too many 'gang' clauses" } */ \
+ seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */ \
+ vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ \
+ worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */
+
+extern void w_4 (void);
+#pragma acc routine (w_4) \
+ worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+ seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */ \
+ vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ \
+ gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */
+#pragma acc routine (w_4) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*w_4." } */ \
+ worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+ vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ \
+ seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */ \
+ gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */
+#pragma acc routine (w_4) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*w_4." } */ \
+ worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+ seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */ \
+ gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */ \
+ vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */
+
+#pragma acc routine \
+ vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+ worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */ \
+ seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */ \
+ gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */
+void v_4 (void)
+{
+}
+#pragma acc routine (v_4) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*v_4." } */ \
+ vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+ seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */ \
+ worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */ \
+ gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */
+#pragma acc routine (v_4) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*v_4." } */ \
+ vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+ gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */ \
+ seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */ \
+ worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */
+
+#pragma acc routine \
+ seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+ vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ \
+ worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */ \
+ gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */
+void s_4 (void)
+{
+}
+#pragma acc routine (s_4) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*s_4." } */ \
+ seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+ vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ \
+ gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */ \
+ worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */
+#pragma acc routine (s_4) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*s_4." } */ \
+ seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+ worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */ \
+ vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ \
+ gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */
+
+
+#pragma acc routine \
+ gang gang gang /* { dg-error "too many 'gang' clauses" } */ \
+ worker worker /* { dg-error "too many 'worker' clauses" } */ \
+ /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+ vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+ /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+ seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+ /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
+void g_5 (void)
+{
+}
+#pragma acc routine (g_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*g_5." } */ \
+ gang gang gang /* { dg-error "too many 'gang' clauses" } */ \
+ vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+ /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+ worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+ /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+ seq seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+ /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
+#pragma acc routine (g_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*g_5." } */ \
+ gang gang gang /* { dg-error "too many 'gang' clauses" } */ \
+ seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+ /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+ vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+ /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+ worker worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+ /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
+
+#pragma acc routine \
+ worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+ seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+ /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+ vector vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+ /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+ gang gang /* { dg-error "too many 'gang' clauses" } */ \
+ /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
+extern void w_5 (void);
+#pragma acc routine (w_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*w_5." } */ \
+ worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+ vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+ /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+ seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+ /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+ gang gang /* { dg-error "too many 'gang' clauses" } */ \
+ /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
+#pragma acc routine (w_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*w_5." } */ \
+ worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+ seq seq /* { dg-error "too many 'seq' clauses" } */ \
+ /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+ gang gang /* { dg-error "too many 'gang' clauses" } */ \
+ /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+ vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+ /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
+
+#pragma acc routine \
+ vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+ worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+ /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+ seq seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+ /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+ gang gang /* { dg-error "too many 'gang' clauses" } */ \
+ /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
+extern void v_5 (void);
+#pragma acc routine (v_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*v_5." } */ \
+ vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+ seq seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+ /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+ worker worker /* { dg-error "too many 'worker' clauses" } */ \
+ /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+ gang gang /* { dg-error "too many 'gang' clauses" } */ \
+ /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
+#pragma acc routine (v_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*v_5." } */ \
+ vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+ gang gang /* { dg-error "too many 'gang' clauses" } */ \
+ /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+ seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+ /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+ worker worker /* { dg-error "too many 'worker' clauses" } */ \
+ /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
+
+extern void s_5 (void);
+#pragma acc routine (s_5) \
+ seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+ vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+ /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+ worker worker worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+ /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+ gang gang /* { dg-error "too many 'gang' clauses" } */ \
+ /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
+#pragma acc routine (s_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*s_5." } */ \
+ seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+ vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+ /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+ gang gang /* { dg-error "too many 'gang' clauses" } */ \
+ /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+ worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+ /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
+#pragma acc routine (s_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*s_5." } */ \
+ seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+ worker worker /* { dg-error "too many 'worker' clauses" } */ \
+ /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+ vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+ /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+ gang gang /* { dg-error "too many 'gang' clauses" } */ \
+ /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
--
2.17.1
>From 33dacef99912c0fb7aa832fb241ac32f7130a1e8 Mon Sep 17 00:00:00 2001
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Fri, 17 May 2019 19:13:26 +0000
Subject: [PATCH 3/3] [PR89433] Repeated use of the C/C++ OpenACC 'routine'
directive
gcc/
PR middle-end/89433
* omp-general.c (oacc_verify_routine_clauses): Change formal
parameters. Add checking if already marked with an OpenACC
'routine' directive. Adjust all users.
gcc/c/
PR c/89433
* c-parser.c (c_finish_oacc_routine): Rework checking if already
marked with an OpenACC 'routine' directive.
gcc/cp/
PR c++/89433
* parser.c (cp_finalize_oacc_routine): Rework checking if already
marked with an OpenACC 'routine' directive.
gcc/testsuite/
PR testsuite/89433
* c-c++-common/goacc/routine-5.c: Update.
* c-c++-common/goacc/routine-level-of-parallelism-1.c: Likewise.
* c-c++-common/goacc/routine-level-of-parallelism-2.c: New file.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@271345 138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/ChangeLog | 5 +
gcc/c/ChangeLog | 4 +
gcc/c/c-parser.c | 46 ++--
gcc/cp/ChangeLog | 4 +
gcc/cp/parser.c | 50 ++--
gcc/omp-general.c | 81 +++++-
gcc/omp-general.h | 3 +-
gcc/testsuite/ChangeLog | 5 +
gcc/testsuite/c-c++-common/goacc/routine-5.c | 46 +---
.../goacc/routine-level-of-parallelism-1.c | 233 ++++++++++++++++--
.../goacc/routine-level-of-parallelism-2.c | 71 ++++++
.../goacc/routine-level-of-parallelism-1.f90 | 6 +-
12 files changed, 437 insertions(+), 117 deletions(-)
create mode 100644 gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-2.c
diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index b8a435860525..bee1292ac229 100644
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,5 +1,10 @@
2019-05-17 Thomas Schwinge <tho...@codesourcery.com>
+ PR middle-end/89433
+ * omp-general.c (oacc_verify_routine_clauses): Change formal
+ parameters. Add checking if already marked with an OpenACC
+ 'routine' directive. Adjust all users.
+
PR middle-end/89433
* omp-general.c (oacc_build_routine_dims): Move some of its
processing into...
diff --git a/gcc/c/ChangeLog b/gcc/c/ChangeLog
index 1393e8f47fd2..cfbd164fdc3c 100644
--- a/gcc/c/ChangeLog
+++ b/gcc/c/ChangeLog
@@ -1,5 +1,9 @@
2019-05-17 Thomas Schwinge <tho...@codesourcery.com>
+ PR c/89433
+ * c-parser.c (c_finish_oacc_routine): Rework checking if already
+ marked with an OpenACC 'routine' directive.
+
PR c/89433
* c-parser.c (c_parser_oacc_routine): Normalize order of clauses.
(c_finish_oacc_routine): Call oacc_verify_routine_clauses.
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 8337f1cce0cf..8f610242435f 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -15884,35 +15884,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/ChangeLog b/gcc/cp/ChangeLog
index 2f1e06ca458f..39aaab966c90 100644
--- a/gcc/cp/ChangeLog
+++ b/gcc/cp/ChangeLog
@@ -1,5 +1,9 @@
2019-05-17 Thomas Schwinge <tho...@codesourcery.com>
+ PR c++/89433
+ * parser.c (cp_finalize_oacc_routine): Rework checking if already
+ marked with an OpenACC 'routine' directive.
+
PR c++/89433
* parser.c (cp_parser_oacc_routine)
(cp_parser_late_parsing_oacc_routine): Normalize order of clauses.
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index aa6507e42f41..6705d64389c1 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -40272,36 +40272,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 f1d859b0275c..82f0a04eab01 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.
+ Returns 0 if FNDECL should be marked with an OpenACC 'routine' directive, 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;
@@ -659,6 +662,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 with an OpenACC 'routine' directive",
+ 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 with an OpenACC 'routine' directive",
+ 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' directive clauses to generate an attribute
diff --git a/gcc/omp-general.h b/gcc/omp-general.h
index 4241c33d99e0..f96d3c7768a0 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/ChangeLog b/gcc/testsuite/ChangeLog
index 4b07888dadbc..ed94f19e1dbd 100644
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,5 +1,10 @@
2019-05-17 Thomas Schwinge <tho...@codesourcery.com>
+ PR testsuite/89433
+ * c-c++-common/goacc/routine-5.c: Update.
+ * c-c++-common/goacc/routine-level-of-parallelism-1.c: Likewise.
+ * c-c++-common/goacc/routine-level-of-parallelism-2.c: New file.
+
PR testsuite/89433
* c-c++-common/goacc/routine-2.c: Update, and move some test
into...
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-5.c b/gcc/testsuite/c-c++-common/goacc/routine-5.c
index b967a7447bdb..a68c6be9be5d 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 with an OpenACC 'routine' directive" } */
#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 with an OpenACC 'routine' directive" } */
#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 ab0414bfed69..4fdeb1461f81 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
@@ -41,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" } */
@@ -52,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" } */
@@ -65,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" } */
@@ -76,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" } */
@@ -90,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" } */ \
@@ -107,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" } */ \
@@ -126,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" } */ \
@@ -145,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" } */ \
@@ -168,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 } */ \
@@ -176,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 } */ \
@@ -194,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 } */ \
@@ -202,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 } */ \
@@ -220,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 } */ \
@@ -228,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 } */ \
@@ -246,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 } */ \
@@ -254,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 } */ \
@@ -262,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 with an OpenACC 'routine' directive" "" { 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 with an OpenACC 'routine' directive" "" { 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 with an OpenACC 'routine' directive" "" { 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 with an OpenACC 'routine' directive" "" { 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 with an OpenACC 'routine' directive" "" { 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 with an OpenACC 'routine' directive" "" { 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 with an OpenACC 'routine' directive" "" { 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 with an OpenACC 'routine' directive" "" { 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 with an OpenACC 'routine' directive" } */
+#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 with an OpenACC 'routine' directive" } */
+
+#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 with an OpenACC 'routine' directive" } */
+#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 with an OpenACC 'routine' directive" } */
+
+#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 with an OpenACC 'routine' directive" } */
+#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 with an OpenACC 'routine' directive" } */
+
+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 with an OpenACC 'routine' directive" } */
+#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 with an OpenACC 'routine' directive" } */
+
+
+/* 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 with an OpenACC 'routine' directive" } */
+
+#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 with an OpenACC 'routine' directive" } */
+
+#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 with an OpenACC 'routine' directive" } */
+
+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 with an OpenACC 'routine' directive" } */
+#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 with an OpenACC 'routine' directive" } */
+#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 with an OpenACC 'routine' directive" } */
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 000000000000..a066f2b9c2ba
--- /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 75dd1b01f6f8..83b8c24b41d8 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
