Hi!
We found that it's not correct that we currently unconditionally diagnose
an error for repeated use of the OpenACC routine directive on one
function/declaration. (For reference, it is also permissible for an
"ordinary" function to have several declarations plus a definition, as
long as these are compatible.) This is, the following shall be valid:
#pragma acc routine worker
void f(void)
{
}
#pragma acc routine (f) worker
#pragma acc routine worker
extern void f(void);
As it is not detailed in the specification, the semantics of repeated use
of OpenACC routine will be implementation-defined. Within one
translation unit, we just remove the existing diagnostic, and declare it
user error if any two usages are not compatible, but we try to be
helpful, and produce a compile-time diagnostic. For incompatible use
spanning over multiple translation units, we're not currently attempting
to produce meaningful diagnostics/semantics; this would be more difficult
(out of scope at present).
Depending on other patches posted (such as "Use
verify_oacc_routine_clauses for C/C++"), here is a patch to implement
this for C/C++ -- Fortran is blocked on the issues raised in the "Use
verify_oacc_routine_clauses for Fortran" patch submission.
One change is to attach an OpenACC routine directive's clauses to the
"omp declare target" attribute (in spirit similar to how they're attached
to OMP_*/GIMPLE_OMP_* codes). (It will then also be possible to get rid
of the "oacc function" attribute and instead have these represented as
OpenACC gang/worker/vector/seq clauses attached to "omp declare target".)
While currently not used for OpenMP, it won't do any harm: continue using
NULL_TREE there, the "empty" clause list. As discussed months ago, on
gomp-4_0-branch we're already using that for implementing the OpenACC
bind and nohost clauses.
Is this conceptually OK for trunk?
commit 91f0742f64f9f0671480023c9949a596691e8476
Author: Thomas Schwinge <tho...@codesourcery.com>
Date: Fri Jul 29 17:47:47 2016 +0200
Repeated use of the OpenACC routine directive
gcc/
* omp-low.c (verify_oacc_routine_clauses): Change formal
parameters. Add checking if already marked as an accelerator
routine. Adjust all users.
gcc/c-family/
* c-common.c (c_common_attribute_table): Set -1 max_len for "omp
declare target".
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/fortran/
* f95-lang.c (gfc_attribute_table): Set -1 max_len for "omp
declare target".
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/c-family/c-common.c | 2 +-
gcc/c/c-parser.c | 42 ++--
gcc/cp/parser.c | 47 +++--
gcc/fortran/f95-lang.c | 2 +-
gcc/fortran/trans-openmp.c | 4 +-
gcc/omp-low.c | 61 +++++-
gcc/omp-low.h | 2 +-
.../c-c++-common/goacc/oaccdevlow-routine.c | 2 +-
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 +++++++
11 files changed, 395 insertions(+), 117 deletions(-)
diff --git gcc/c-family/c-common.c gcc/c-family/c-common.c
index 16e3965..63ed75a 100644
--- gcc/c-family/c-common.c
+++ gcc/c-family/c-common.c
@@ -822,7 +822,7 @@ const struct attribute_spec c_common_attribute_table[] =
handle_omp_declare_simd_attribute, false },
{ "simd", 0, 1, true, false, false,
handle_simd_attribute, false },
- { "omp declare target", 0, 0, true, false, false,
+ { "omp declare target", 0, -1, true, false, false,
handle_omp_declare_target_attribute, false },
{ "omp declare target link", 0, 0, true, false, false,
handle_omp_declare_target_attribute, false },
diff --git gcc/c/c-parser.c gcc/c/c-parser.c
index 24f60cf..8e67ed8 100644
--- gcc/c/c-parser.c
+++ gcc/c/c-parser.c
@@ -14155,33 +14155,37 @@ c_finish_oacc_routine (struct oacc_routine_data
*data, tree fndecl,
return;
}
- verify_oacc_routine_clauses (&data->clauses, data->loc);
-
- if (get_oacc_fn_attrib (fndecl))
+ int compatible
+ = verify_oacc_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,
- "%<#pragma acc routine%> must be applied before %s",
- TREE_USED (fndecl) ? "use" : "definition");
- data->error_seen = true;
- return;
}
+ else
+ {
+ if (TREE_USED (fndecl) || (!is_defn && DECL_SAVED_TREE (fndecl)))
+ {
+ error_at (data->loc,
+ "%<#pragma acc routine%> must be applied before %s",
+ TREE_USED (fndecl) ? "use" : "definition");
+ data->error_seen = true;
+ return;
+ }
- /* Process the routine's dimension clauses. */
- tree dims = build_oacc_routine_dims (data->clauses);
- replace_oacc_fn_attrib (fndecl, dims);
+ /* Set the routine's level of parallelism. */
+ tree dims = build_oacc_routine_dims (data->clauses);
+ replace_oacc_fn_attrib (fndecl, dims);
- /* Add an "omp declare target" attribute. */
- DECL_ATTRIBUTES (fndecl)
- = tree_cons (get_identifier ("omp declare target"),
- NULL_TREE, 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 gcc/cp/parser.c gcc/cp/parser.c
index 6197fc9..34afa17 100644
--- gcc/cp/parser.c
+++ gcc/cp/parser.c
@@ -36750,34 +36750,39 @@ cp_finalize_oacc_routine (cp_parser *parser, tree
fndecl, bool is_defn)
return;
}
- verify_oacc_routine_clauses (&parser->oacc_routine->clauses,
- parser->oacc_routine->loc);
-
- if (get_oacc_fn_attrib (fndecl))
+ int compatible
+ = verify_oacc_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,
- "%<#pragma acc routine%> must be applied before %s",
- TREE_USED (fndecl) ? "use" : "definition");
- parser->oacc_routine = NULL;
- return;
}
+ else
+ {
+ if (TREE_USED (fndecl) || (!is_defn && DECL_SAVED_TREE (fndecl)))
+ {
+ error_at (parser->oacc_routine->loc,
+ "%<#pragma acc routine%> must be applied before %s",
+ TREE_USED (fndecl) ? "use" : "definition");
+ parser->oacc_routine = NULL;
+ return;
+ }
+
+ /* Set the routine's level of parallelism. */
+ tree dims = build_oacc_routine_dims (parser->oacc_routine->clauses);
+ replace_oacc_fn_attrib (fndecl, dims);
- /* Process the routine's dimension clauses. */
- tree dims = build_oacc_routine_dims (parser->oacc_routine->clauses);
- replace_oacc_fn_attrib (fndecl, dims);
-
- /* Add an "omp declare target" attribute. */
- DECL_ATTRIBUTES (fndecl)
- = tree_cons (get_identifier ("omp declare target"),
- NULL_TREE, 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 gcc/fortran/f95-lang.c gcc/fortran/f95-lang.c
index 2b58173..5f6e54f 100644
--- gcc/fortran/f95-lang.c
+++ gcc/fortran/f95-lang.c
@@ -90,7 +90,7 @@ static const struct attribute_spec gfc_attribute_table[] =
{
/* { name, min_len, max_len, decl_req, type_req, fn_type_req, handler,
affects_type_identity } */
- { "omp declare target", 0, 0, true, false, false,
+ { "omp declare target", 0, -1, true, false, false,
gfc_handle_omp_declare_target_attribute, false },
{ "oacc function", 0, -1, true, false, false,
gfc_handle_omp_declare_target_attribute, false },
diff --git gcc/fortran/trans-openmp.c gcc/fortran/trans-openmp.c
index 254732c..0239871 100644
--- gcc/fortran/trans-openmp.c
+++ gcc/fortran/trans-openmp.c
@@ -4612,7 +4612,9 @@ gfc_oacc_routine_dims (gfc_omp_clauses *clauses, locus
location)
clauses_t = c;
}
}
- verify_oacc_routine_clauses (&clauses_t, loc);
+ (void) verify_oacc_routine_clauses (/* TODO */ NULL_TREE,
+ &clauses_t, loc,
+ "!$ACC ROUTINE");
gcc_checking_assert (clauses_t != NULL_TREE
&& OMP_CLAUSE_CHAIN (clauses_t) == NULL_TREE);
diff --git gcc/omp-low.c gcc/omp-low.c
index 74b8a99..8a8f222 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -12597,11 +12597,14 @@ set_oacc_fn_attrib (tree fn, tree clauses, vec<tree>
*args)
/* Verify OpenACC routine clauses.
- Upon returning, the chain of clauses 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
-verify_oacc_routine_clauses (tree *clauses, location_t loc)
+int
+verify_oacc_routine_clauses (tree fndecl, tree *clauses, location_t loc,
+ const char *routine_str)
{
tree c_level = NULL_TREE;
tree c_p = NULL_TREE;
@@ -12647,6 +12650,56 @@ verify_oacc_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. */
+
+ /* Still got some work to do for Fortran... */
+ if (fndecl == NULL_TREE)
+ return 0;
+
+ 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. */
+ if (OMP_CLAUSE_CODE (c_level) != OMP_CLAUSE_CODE (c_level_p))
+ {
+ error_at (OMP_CLAUSE_LOCATION (c_level),
+ "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_level)],
+ routine_str, fndecl);
+ inform (OMP_CLAUSE_LOCATION (c_level_p),
+ "... with %qs clause here",
+ omp_clause_code_name[OMP_CLAUSE_CODE (c_level_p)]);
+ /* Incompatible. */
+ return -1;
+ }
+ /* Compatible. */
+ return 1;
+ }
+
+ return 0;
}
/* Process the OpenACC routine's clauses to generate an attribute
diff --git gcc/omp-low.h gcc/omp-low.h
index f02d990..0ea5308 100644
--- gcc/omp-low.h
+++ gcc/omp-low.h
@@ -31,7 +31,7 @@ extern bool make_gimple_omp_edges (basic_block, struct
omp_region **, int *);
extern void omp_finish_file (void);
extern tree omp_member_access_dummy_var (tree);
extern void replace_oacc_fn_attrib (tree, tree);
-extern void verify_oacc_routine_clauses (tree *, location_t);
+extern int verify_oacc_routine_clauses (tree, tree *, location_t, const char
*);
extern tree build_oacc_routine_dims (tree);
extern tree get_oacc_fn_attrib (tree);
extern void set_oacc_fn_attrib (tree, tree, vec<tree> *);
diff --git gcc/testsuite/c-c++-common/goacc/oaccdevlow-routine.c
gcc/testsuite/c-c++-common/goacc/oaccdevlow-routine.c
index bb02fc7..771f3b1 100644
--- gcc/testsuite/c-c++-common/goacc/oaccdevlow-routine.c
+++ gcc/testsuite/c-c++-common/goacc/oaccdevlow-routine.c
@@ -21,7 +21,7 @@ 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, 1, 1] for target compilation).
diff --git gcc/testsuite/c-c++-common/goacc/routine-5.c
gcc/testsuite/c-c++-common/goacc/routine-5.c
index 17fe67c..f10651d 100644
--- gcc/testsuite/c-c++-common/goacc/routine-5.c
+++ 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 gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-1.c
gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-1.c
index fc8bc3c..8f45499 100644
--- gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-1.c
+++ gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-1.c
@@ -42,10 +42,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 +53,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 +66,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 +77,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 +91,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 +108,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 +127,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 +146,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 +169,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 *-*-* } 174 } */ \
@@ -177,7 +177,7 @@ void g_5 (void)
/* { dg-error ".worker. specifies a conflicting level of parallelism" "" {
target *-*-* } 176 } */ \
seq seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
/* { dg-error ".seq. specifies a conflicting level of parallelism" "" {
target *-*-* } 178 } */
-#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 *-*-* } 182 } */ \
@@ -195,7 +195,7 @@ void g_5 (void)
gang gang /* { dg-error "too many 'gang' clauses" } */ \
/* { dg-error ".gang. specifies a conflicting level of parallelism" "" {
target *-*-* } 195 } */
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 *-*-* } 200 } */ \
@@ -203,7 +203,7 @@ extern void w_5 (void);
/* { dg-error ".seq. specifies a conflicting level of parallelism" "" {
target *-*-* } 202 } */ \
gang gang /* { dg-error "too many 'gang' clauses" } */ \
/* { dg-error ".gang. specifies a conflicting level of parallelism" "" {
target *-*-* } 204 } */
-#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 *-*-* } 208 } */ \
@@ -221,7 +221,7 @@ extern void w_5 (void);
gang gang /* { dg-error "too many 'gang' clauses" } */ \
/* { dg-error ".gang. specifies a conflicting level of parallelism" "" {
target *-*-* } 221 } */
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 *-*-* } 226 } */ \
@@ -229,7 +229,7 @@ extern void v_5 (void);
/* { dg-error ".worker. specifies a conflicting level of parallelism" "" {
target *-*-* } 228 } */ \
gang gang /* { dg-error "too many 'gang' clauses" } */ \
/* { dg-error ".gang. specifies a conflicting level of parallelism" "" {
target *-*-* } 230 } */
-#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 *-*-* } 234 } */ \
@@ -247,7 +247,7 @@ extern void s_5 (void);
/* { dg-error ".worker. specifies a conflicting level of parallelism" "" {
target *-*-* } 246 } */ \
gang gang /* { dg-error "too many 'gang' clauses" } */ \
/* { dg-error ".gang. specifies a conflicting level of parallelism" "" {
target *-*-* } 248 } */
-#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 *-*-* } 252 } */ \
@@ -255,7 +255,7 @@ extern void s_5 (void);
/* { dg-error ".gang. specifies a conflicting level of parallelism" "" {
target *-*-* } 254 } */ \
worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
/* { dg-error ".worker. specifies a conflicting level of parallelism" "" {
target *-*-* } 256 } */
-#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 *-*-* } 260 } */ \
@@ -263,3 +263,188 @@ extern void s_5 (void);
/* { dg-error ".vector. specifies a conflicting level of parallelism" "" {
target *-*-* } 262 } */ \
gang gang /* { dg-error "too many 'gang' clauses" } */ \
/* { dg-error ".gang. specifies a conflicting level of parallelism" "" {
target *-*-* } 264 } */
+
+
+/* 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 *-*-* } 273 } */ \
+ vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+ /* { dg-error ".vector. specifies a conflicting level of parallelism" "" {
target *-*-* } 275 } */ \
+ seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+ /* { dg-error ".seq. specifies a conflicting level of parallelism" "" {
target *-*-* } 277 } */
+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 *-*-* } 283 } */ \
+ gang gang gang /* { dg-error "too many 'gang' clauses" } */ \
+ /* { dg-error ".gang. specifies a conflicting level of parallelism" "" {
target *-*-* } 285 } */ \
+ worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+ /* { dg-error ".worker. specifies a conflicting level of parallelism" "" {
target *-*-* } 287 } */ \
+ seq seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+ /* { dg-error ".seq. specifies a conflicting level of parallelism" "" {
target *-*-* } 289 } */
+#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 *-*-* } 292 } */ \
+ gang gang gang /* { dg-error "too many 'gang' clauses" } */ \
+ /* { dg-error ".gang. specifies a conflicting level of parallelism" "" {
target *-*-* } 294 } */ \
+ vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+ /* { dg-error ".vector. specifies a conflicting level of parallelism" "" {
target *-*-* } 296 } */ \
+ worker worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+ /* { dg-error ".worker. specifies a conflicting level of parallelism" "" {
target *-*-* } 298 } */
+
+#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 *-*-* } 303 } */ \
+ vector vector vector vector vector /* { dg-error "too many 'vector' clauses"
} */ \
+ /* { dg-error ".vector. specifies a conflicting level of parallelism" "" {
target *-*-* } 305 } */ \
+ gang gang /* { dg-error "too many 'gang' clauses" } */ \
+ /* { dg-error ".gang. specifies a conflicting level of parallelism" "" {
target *-*-* } 307 } */
+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 *-*-* } 311 } */ \
+ worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+ /* { dg-error ".worker. specifies a conflicting level of parallelism" "" {
target *-*-* } 313 } */ \
+ seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+ /* { dg-error ".seq. specifies a conflicting level of parallelism" "" {
target *-*-* } 315 } */ \
+ gang gang /* { dg-error "too many 'gang' clauses" } */ \
+ /* { dg-error ".gang. specifies a conflicting level of parallelism" "" {
target *-*-* } 317 } */
+#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 *-*-* } 320 } */ \
+ worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+ /* { dg-error ".worker. specifies a conflicting level of parallelism" "" {
target *-*-* } 322 } */ \
+ gang gang /* { dg-error "too many 'gang' clauses" } */ \
+ /* { dg-error ".gang. specifies a conflicting level of parallelism" "" {
target *-*-* } 324 } */ \
+ vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+ /* { dg-error ".vector. specifies a conflicting level of parallelism" "" {
target *-*-* } 326 } */
+
+#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 *-*-* } 331 } */ \
+ seq seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+ /* { dg-error ".seq. specifies a conflicting level of parallelism" "" {
target *-*-* } 333 } */ \
+ gang gang /* { dg-error "too many 'gang' clauses" } */ \
+ /* { dg-error ".gang. specifies a conflicting level of parallelism" "" {
target *-*-* } 335 } */
+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 *-*-* } 339 } */ \
+ vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+ /* { dg-error ".vector. specifies a conflicting level of parallelism" "" {
target *-*-* } 341 } */ \
+ worker worker /* { dg-error "too many 'worker' clauses" } */ \
+ /* { dg-error ".worker. specifies a conflicting level of parallelism" "" {
target *-*-* } 343 } */ \
+ gang gang /* { dg-error "too many 'gang' clauses" } */ \
+ /* { dg-error ".gang. specifies a conflicting level of parallelism" "" {
target *-*-* } 345 } */
+#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 *-*-* } 348 } */ \
+ vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+ /* { dg-error ".vector. specifies a conflicting level of parallelism" "" {
target *-*-* } 350 } */ \
+ seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+ /* { dg-error ".seq. specifies a conflicting level of parallelism" "" {
target *-*-* } 352 } */ \
+ worker worker /* { dg-error "too many 'worker' clauses" } */ \
+ /* { dg-error ".worker. specifies a conflicting level of parallelism" "" {
target *-*-* } 354 } */
+
+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 *-*-* } 360 } */ \
+ worker worker worker worker worker /* { dg-error "too many 'worker' clauses"
} */ \
+ /* { dg-error ".worker. specifies a conflicting level of parallelism" "" {
target *-*-* } 362 } */ \
+ gang gang /* { dg-error "too many 'gang' clauses" } */ \
+ /* { dg-error ".gang. specifies a conflicting level of parallelism" "" {
target *-*-* } 364 } */
+#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 *-*-* } 367 } */ \
+ seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+ /* { dg-error ".seq. specifies a conflicting level of parallelism" "" {
target *-*-* } 369 } */ \
+ gang gang /* { dg-error "too many 'gang' clauses" } */ \
+ /* { dg-error ".gang. specifies a conflicting level of parallelism" "" {
target *-*-* } 371 } */ \
+ worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+ /* { dg-error ".worker. specifies a conflicting level of parallelism" "" {
target *-*-* } 373 } */
+#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 *-*-* } 376 } */ \
+ seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+ /* { dg-error ".seq. specifies a conflicting level of parallelism" "" {
target *-*-* } 378 } */ \
+ vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+ /* { dg-error ".vector. specifies a conflicting level of parallelism" "" {
target *-*-* } 380 } */ \
+ gang gang /* { dg-error "too many 'gang' clauses" } */ \
+ /* { dg-error ".gang. specifies a conflicting level of parallelism" "" {
target *-*-* } 382 } */
+
+
+/* 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 gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-2.c
gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-2.c
new file mode 100644
index 0000000..1217397
--- /dev/null
+++ 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-2.f. */
+
+#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
Grüße
Thomas
