This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
[gomp4] Middle end bits for routines
- From: Nathan Sidwell <nathan at acm dot org>
- To: GCC Patches <gcc-patches at gcc dot gnu dot org>
- Date: Sat, 15 Aug 2015 11:21:09 -0400
- Subject: [gomp4] Middle end bits for routines
- Authentication-results: sourceware.org; auth=none
I've committed this to the gomp4 branch. It extends the 'oacc function'
attribute's dimension handling to deal with routines. With the latter we now
use TREE_PURPOSE of the dimension list to indicate whether the routine may or
may not spawn a partitioned loop on a particular axis. I had to tweak the
validate_dims hook to tell it what the outermost such axis is (or -1 for a
non-routine).
The patch was complicated y the rather baroque handling of the routine directive
in the C and C++ frontends. Now I see that handling is based on a
misinterpretation of the specification. I'll clean that up shortly.
nathan
2015-08-15 Nathan Sidwell <nathan@codesourcery.com>
* config/nvptx/nvptx.c (nvptx_reorg): Examine TREE_PURPOSE of
dimensions.
(nvptx_record_offload_symbol): Adjust.
(nvptx_validate_dims): Adjust.
* omp-low.c (replace_oacc_fn_attrib): Detect if we're the first
attribute.
(set_oacc_fn_attrub): Swap FN and CLAUSE parameters for
consistency.
(build_oacc_routine_dims): New.
(expand_omp_target): Adjust set_oacc_fn_attrib call.
(execute_oacc_transform): Deal with routine dimensions.
(default_goacc_validate_dims): Add FN_LEVEL parameter.
* omp-low.h (replace_oacc_fn_attrib, build_oacc_routine_dims): Declare.
* target.def (validate_dims): Add FN_LEVEL arg.
* targhooks.h (default_goacc_validate_dims): Adjust.
* doc/tm.texi: Rebuilt.
testsuite/
* c-c++-common/goacc/routine-4.c: Adjust expected error. Delete
bogus tests.
fortran/
* f95-lang.c (gfc_attribs): oacc function attribute can take
list.
* gfortran.h (struct symbol_attribute): Add more bits to
oacc_routine field.
* openmp.c (gfc_oacc_routine_dims): New.
(gfc_match_oacc_routins): Call it.
* trans-decl.c: Include gomp-constants.h.
(add_attributes_to_decl): Create oacc function dimension data.
cp/
* parser.c (cp_parser_oacc_routine_check_parallelism): Delete.
(cp_parser_oacc_routine): Don't check parallelism here.
(cp_parser_late_parssing_oacc_routine): Use build_oacc_routine_dims.
c/
* c-parser.c (c_parser_oacc_routine): Don't check loop axes here,
use build_oacc_routine_dims.
c-family/
* c-common.c (c_common_attribs): oacc function can take list.
Index: gcc/c/c-parser.c
===================================================================
--- gcc/c/c-parser.c (revision 226860)
+++ gcc/c/c-parser.c (working copy)
@@ -13291,7 +13291,6 @@ static void
c_parser_oacc_routine (c_parser *parser, enum pragma_context context)
{
tree name = NULL_TREE;
- location_t here = c_parser_peek_token (parser)->location;
c_parser_consume_pragma (parser);
@@ -13329,30 +13328,6 @@ c_parser_oacc_routine (c_parser *parser,
"#pragma acc routine",
OACC_ROUTINE_CLAUSE_DEVICE_TYPE_MASK);
- /* Check of the presence if gang, worker, vector and seq clauses, and
- throw an error if more than one of those clauses is specified. */
- int parallelism = 0;
- tree c;
-
- for (c = clauses; 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:
- ++parallelism;
- break;
- default:
- break;
- }
-
- if (parallelism > 1)
- {
- error_at (here, "invalid combination of gang, worker, vector or seq for"
- "%<#pragma acc routine%>");
- }
-
if (name)
{
TREE_CHAIN (name) = clauses;
@@ -13422,14 +13397,16 @@ c_finish_oacc_routine (c_parser *parser,
return;
}
+ /* Process for function attrib */
+ tree dims = build_oacc_routine_dims (clauses);
+ replace_oacc_fn_attrib (fndecl, dims);
+
+ /* Also attach as a declare. */
if (clauses != NULL_TREE)
clauses = tree_cons (NULL_TREE, clauses, NULL_TREE);
- clauses = build_tree_list (get_identifier ("omp declare target"),
- clauses);
- TREE_CHAIN (clauses) = DECL_ATTRIBUTES (fndecl);
DECL_ATTRIBUTES (fndecl)
- = tree_cons (get_identifier ("oacc function"),
- NULL_TREE, clauses);
+ = tree_cons (get_identifier ("omp declare target"),
+ clauses, DECL_ATTRIBUTES (fndecl));
}
/* OpenACC 2.0:
Index: gcc/c-family/c-common.c
===================================================================
--- gcc/c-family/c-common.c (revision 226860)
+++ gcc/c-family/c-common.c (working copy)
@@ -823,7 +823,7 @@ const struct attribute_spec c_common_att
{ "bnd_instrument", 0, 0, true, false, false,
handle_bnd_instrument, false },
{ "oacc declare", 0, -1, true, false, false, NULL, false },
- { "oacc function", 0, 0, true, false, false, NULL, false },
+ { "oacc function", 0, -1, true, false, false, NULL, false },
{ NULL, 0, 0, false, false, false, NULL, false }
};
Index: gcc/config/nvptx/nvptx.c
===================================================================
--- gcc/config/nvptx/nvptx.c (revision 226860)
+++ gcc/config/nvptx/nvptx.c (working copy)
@@ -3079,9 +3079,10 @@ nvptx_reorg (void)
for (ix = 0; ix != GOMP_DIM_MAX; ix++, dims = TREE_CHAIN (dims))
{
- HOST_WIDE_INT size = TREE_INT_CST_LOW (TREE_VALUE (dims));
+ int size = TREE_INT_CST_LOW (TREE_VALUE (dims));
+ tree allowed = TREE_PURPOSE (dims);
- if (size > 1 || (!size && !TREE_PURPOSE (dims)))
+ if (size != 1 && !(allowed && integer_zerop (allowed)))
mask |= GOMP_DIM_MASK (ix);
}
/* If there is worker neutering, there must be vector
@@ -3188,10 +3189,10 @@ nvptx_record_offload_symbol (tree decl)
for (ix = 0; ix != GOMP_DIM_MAX; ix++, dims = TREE_CHAIN (dims))
{
- HOST_WIDE_INT size = TREE_INT_CST_LOW (TREE_VALUE (dims));
+ int size = TREE_INT_CST_LOW (TREE_VALUE (dims));
gcc_assert (!TREE_PURPOSE (dims));
- fprintf (asm_out_file, ", " HOST_WIDE_INT_PRINT_HEX, size);
+ fprintf (asm_out_file, ", %#x", size);
}
fprintf (asm_out_file, "\n");
@@ -3543,15 +3544,24 @@ nvptx_expand_builtin (tree exp, rtx targ
#define PTX_VECTOR_LENGTH 32
#define PTX_WORKER_LENGTH 32
-/* Validate compute dimensions, fill in non-unity defaults. */
+/* Validate compute dimensions, fill in non-unity defaults. FN_LEVEL
+ indicates the level at which a routine might spawn a loop. It is
+ negative for non-routines. */
static bool
-nvptx_validate_dims (tree decl, int dims[])
+nvptx_validate_dims (tree decl, int dims[], int fn_level)
{
bool changed = false;
+ if (fn_level >= 0)
+ /* This is a routine. All dimensions are dynamic and controlled
+ by the calling function. Because we permit a 1vx1wxNg
+ geometry, we can't take the opportunity to fix the vector
+ dimension inside a routine. Perhaps we should? */
+ return false;
+
/* If the worker size is not 1, the vector size must be 32. If
- the vector size is not 1, it must be 32. */
+ the vector size is not 1, it must be 32. */
if ((dims[GOMP_DIM_WORKER] > 1 || dims[GOMP_DIM_WORKER] == 0)
|| (dims[GOMP_DIM_VECTOR] > 1 || dims[GOMP_DIM_VECTOR] == 0))
{
Index: gcc/cp/parser.c
===================================================================
--- gcc/cp/parser.c (revision 226860)
+++ gcc/cp/parser.c (working copy)
@@ -34368,34 +34368,6 @@ cp_parser_omp_declare (cp_parser *parser
cp_parser_require_pragma_eol (parser, pragma_tok);
}
-static void
-cp_parser_oacc_routine_check_parallelism (tree clauses, location_t loc)
-{
- /* Check of the presence if gang, worker, vector and seq clauses, and
- throw an error if more than one of those clauses is specified. */
- int parallelism = 0;
- tree c;
-
- for (c = clauses; 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:
- ++parallelism;
- break;
- default:
- break;
- }
-
- if (parallelism > 1)
- {
- error_at (loc, "invalid combination of gang, worker, vector or seq for"
- "%<#pragma acc routine%>");
- }
-}
-
/* OpenACC 2.0:
# pragma acc routine oacc-routine-clause[optseq] new-line
function-definition
@@ -34424,7 +34396,6 @@ cp_parser_oacc_routine (cp_parser *parse
enum pragma_context context)
{
tree name = NULL_TREE;
- location_t here = cp_lexer_peek_token (parser->lexer)->location;
//cp_lexer_consume_token (parser->lexer);
@@ -34466,8 +34437,6 @@ cp_parser_oacc_routine (cp_parser *parse
cp_lexer_peek_token (parser->lexer),
OACC_ROUTINE_CLAUSE_DEVICE_TYPE_MASK);
- cp_parser_oacc_routine_check_parallelism (clauses, here);
-
TREE_CHAIN (name) = clauses;
vec_safe_push (parser->named_oacc_routines, name);
}
@@ -34481,7 +34450,6 @@ cp_parser_late_parsing_oacc_routine (cp_
struct cp_token_cache *ce;
cp_omp_declare_simd_data *data = parser->oacc_routine;
int i;
- location_t here = UNKNOWN_LOCATION;
if (!data->error_seen && data->fndecl_seen)
{
@@ -34499,7 +34467,6 @@ cp_parser_late_parsing_oacc_routine (cp_
{
cp_parser_push_lexer_for_tokens (parser, ce);
parser->lexer->in_pragma = true;
- here = cp_lexer_peek_token (parser->lexer)->location;
gcc_assert (cp_lexer_peek_token (parser->lexer)->type == CPP_PRAGMA);
cp_token *pragma_tok = cp_lexer_consume_token (parser->lexer);
c = cp_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK,
@@ -34519,13 +34486,13 @@ cp_parser_late_parsing_oacc_routine (cp_
}
}
- cp_parser_oacc_routine_check_parallelism (cl, here);
+ tree dims = build_oacc_routine_dims (cl);
+ attrs = tree_cons (get_identifier ("oacc function"), dims, attrs);
if (cl != NULL_TREE)
cl = tree_cons (NULL_TREE, cl, NULL_TREE);
- attrs = build_tree_list (get_identifier ("omp declare target"), cl);
- attrs = tree_cons (get_identifier ("oacc function"), NULL_TREE, attrs);
+ attrs = tree_cons (get_identifier ("omp declare target"), cl, attrs);
data->fndecl_seen = true;
return attrs;
}
Index: gcc/doc/tm.texi
===================================================================
--- gcc/doc/tm.texi (revision 226860)
+++ gcc/doc/tm.texi (working copy)
@@ -5740,7 +5740,7 @@ usable. In that case, the smaller the n
to use it.
@end deftypefn
-@deftypefn {Target Hook} bool TARGET_GOACC_VALIDATE_DIMS (tree, int @var{[]})
+@deftypefn {Target Hook} bool TARGET_GOACC_VALIDATE_DIMS (tree, int @var{[]}, @var{int})
This hook should check the launch dimensions provided. It should fill
in anything that needs default to non-unity and verify non-defaults.
Defaults are represented as -1. Diagnostics should be issuedas
Index: gcc/fortran/f95-lang.c
===================================================================
--- gcc/fortran/f95-lang.c (revision 226860)
+++ gcc/fortran/f95-lang.c (working copy)
@@ -106,7 +106,7 @@ static const struct attribute_spec gfc_a
gfc_handle_omp_declare_target_attribute, false },
{ "oacc declare", 0, 0, true, false, false,
gfc_handle_omp_declare_target_attribute, false },
- { "oacc function", 0, 0, true, false, false,
+ { "oacc function", 0, -1, true, false, false,
gfc_handle_omp_declare_target_attribute, false },
{ NULL, 0, 0, false, false, false, NULL, false }
};
Index: gcc/fortran/gfortran.h
===================================================================
--- gcc/fortran/gfortran.h (revision 226860)
+++ gcc/fortran/gfortran.h (working copy)
@@ -875,8 +875,8 @@ typedef struct
unsigned oacc_declare_device_resident:1;
unsigned oacc_declare_link:1;
- /* This is an OpenACC acclerator function. */
- unsigned oacc_function:1;
+ /* This is an OpenACC acclerator function at level N - 1 */
+ unsigned oacc_function:3;
/* Attributes set by compiler extensions (!GCC$ ATTRIBUTES). */
unsigned ext_attr:EXT_ATTR_NUM;
Index: gcc/fortran/openmp.c
===================================================================
--- gcc/fortran/openmp.c (revision 226860)
+++ gcc/fortran/openmp.c (working copy)
@@ -1691,6 +1691,35 @@ gfc_match_oacc_cache (void)
return MATCH_YES;
}
+/* Determine the loop level for a routine. */
+
+static int
+gfc_oacc_routine_dims (gfc_omp_clauses *clauses)
+{
+ int level = -1;
+
+ if (clauses)
+ {
+ unsigned mask = 0;
+
+ if (clauses->gang)
+ level = GOMP_DIM_GANG, mask |= GOMP_DIM_MASK (level);
+ if (clauses->worker)
+ level = GOMP_DIM_WORKER, mask |= GOMP_DIM_MASK (level);
+ if (clauses->vector)
+ level = GOMP_DIM_VECTOR, mask |= GOMP_DIM_MASK (level);
+ if (clauses->seq)
+ level = GOMP_DIM_MAX, mask |= GOMP_DIM_MASK (level);
+
+ if (mask != (mask & -mask))
+ gfc_error ("Multiple loop axes specified for routine");
+ }
+
+ if (level < 0)
+ level = GOMP_DIM_MAX;
+
+ return level;
+}
match
gfc_match_oacc_routine (void)
@@ -1743,6 +1772,12 @@ gfc_match_oacc_routine (void)
}
}
+ if (gfc_match_omp_eos () != MATCH_YES
+ && gfc_match_omp_clauses (&c, OACC_ROUTINE_CLAUSES,
+ OACC_ROUTINE_CLAUSE_DEVICE_TYPE_MASK, false,
+ false, true) != MATCH_YES)
+ return MATCH_ERROR;
+
if (sym != NULL)
{
n = gfc_get_oacc_routine_name ();
@@ -1760,20 +1795,12 @@ gfc_match_oacc_routine (void)
gfc_current_ns->proc_name->name,
&old_loc))
goto cleanup;
- gfc_current_ns->proc_name->attr.oacc_function = 1;
+ gfc_current_ns->proc_name->attr.oacc_function
+ = gfc_oacc_routine_dims (c) + 1;
}
else
gcc_unreachable ();
- if (gfc_match_omp_eos () == MATCH_YES)
- return MATCH_YES;
-
- if (gfc_match_omp_clauses (&c, OACC_ROUTINE_CLAUSES,
- OACC_ROUTINE_CLAUSE_DEVICE_TYPE_MASK, false,
- false, true)
- != MATCH_YES)
- return MATCH_ERROR;
-
if (n)
n->clauses = c;
else if (gfc_current_ns->oacc_routine)
Index: gcc/fortran/trans-decl.c
===================================================================
--- gcc/fortran/trans-decl.c (revision 226860)
+++ gcc/fortran/trans-decl.c (working copy)
@@ -49,6 +49,7 @@ along with GCC; see the file COPYING3.
#include "trans-const.h"
/* Only for gfc_trans_code. Shouldn't need to include this. */
#include "trans-stmt.h"
+#include "gomp-constants.h"
#define MAX_LABEL_VALUE 99999
@@ -1320,8 +1321,18 @@ add_attributes_to_decl (symbol_attribute
}
if (sym_attr.oacc_function)
- list = tree_cons (get_identifier ("oacc function"),
- NULL_TREE, list);
+ {
+ tree dims = NULL_TREE;
+ int ix;
+ int level = sym_attr.oacc_function - 1;
+
+ for (ix = GOMP_DIM_MAX; ix--;)
+ dims = tree_cons (build_int_cst (boolean_type_node, ix >= level),
+ integer_zero_node, dims);
+
+ list = tree_cons (get_identifier ("oacc function"),
+ dims, list);
+ }
return list;
}
Index: gcc/omp-low.c
===================================================================
--- gcc/omp-low.c (revision 226860)
+++ gcc/omp-low.c (working copy)
@@ -9319,13 +9319,17 @@ oacc_launch_pack (unsigned code, tree de
void
replace_oacc_fn_attrib (tree fn, tree dims)
{
- /* Simply cons onto the beginning of the list. */
- DECL_ATTRIBUTES (fn) =
- tree_cons (get_identifier (OACC_FN_ATTRIB), dims, DECL_ATTRIBUTES (fn));
+ tree ident = get_identifier (OACC_FN_ATTRIB);
+ tree attribs = DECL_ATTRIBUTES (fn);
+
+ /* If we happen to be present as the first attrib, drop it. */
+ if (attribs && TREE_PURPOSE (attribs) == ident)
+ attribs = TREE_CHAIN (attribs);
+ DECL_ATTRIBUTES (fn) = tree_cons (ident, dims, attribs);
}
static void
-set_oacc_fn_attrib (tree clauses, tree fn, vec<tree> *args)
+set_oacc_fn_attrib (tree fn, tree clauses, vec<tree> *args)
{
/* Must match GOMP_DIM ordering. */
static const omp_clause_code ids[] =
@@ -9364,6 +9368,45 @@ set_oacc_fn_attrib (tree clauses, tree f
}
}
+/* 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
+ (dynamic). TREE_PURPOSE is set to indicate whether that dimension
+ can have a loop partitioned on it. boolean_true_node indicates
+ yes, boolean_false_node indicates no. */
+
+tree
+build_oacc_routine_dims (tree clauses)
+{
+ /* Must match GOMP_DIM ordering. */
+ static const omp_clause_code ids[] =
+ {OMP_CLAUSE_GANG, OMP_CLAUSE_WORKER, OMP_CLAUSE_VECTOR, OMP_CLAUSE_SEQ};
+ int ix;
+ int level = -1;
+
+ for (; clauses; clauses = OMP_CLAUSE_CHAIN (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;
+ }
+
+ if (level < 0)
+ level = GOMP_DIM_MAX;
+
+ tree dims = NULL_TREE;
+
+ for (ix = GOMP_DIM_MAX; ix--;)
+ dims = tree_cons (build_int_cst (boolean_type_node, ix >= level),
+ integer_zero_node, dims);
+
+ return dims;
+}
+
tree
get_oacc_fn_attrib (tree fn)
{
@@ -9862,7 +9905,7 @@ expand_omp_target (struct omp_region *re
case BUILT_IN_GOACC_PARALLEL:
{
args.quick_push (gimple_omp_target_ganglocal_size (entry_stmt));
- set_oacc_fn_attrib (clauses, child_fn, &args);
+ set_oacc_fn_attrib (child_fn, clauses, &args);
tagging = true;
}
/* FALLTHRU */
@@ -14624,6 +14667,7 @@ execute_oacc_transform ()
basic_block bb;
tree attrs = get_oacc_fn_attrib (current_function_decl);
int dims[GOMP_DIM_MAX];
+ tree purpose[GOMP_DIM_MAX];
if (!attrs)
/* Not an offloaded function. */
@@ -14632,36 +14676,47 @@ execute_oacc_transform ()
{
unsigned ix;
tree pos = TREE_VALUE (attrs);
+ int fn_level = -1;
+ /* Make sure the attribute creator attached the dimension
+ information. */
+ gcc_assert (pos);
+
for (ix = 0; ix != GOMP_DIM_MAX; ix++)
{
- if (!pos)
- dims[ix] = -1;
- else
+ purpose[ix] = TREE_PURPOSE (pos);
+
+ if (purpose[ix])
{
- tree val = TREE_VALUE (pos);
-
- dims[ix] = val ? TREE_INT_CST_LOW (val) : -2;
- pos = TREE_CHAIN (pos);
+ if (purpose[ix] == boolean_false_node)
+ fn_level = ix + 1;
+ else if (fn_level < 0)
+ fn_level = ix;
}
+
+ tree val = TREE_VALUE (pos);
+
+ dims[ix] = val ? TREE_INT_CST_LOW (val) : -1;
+ pos = TREE_CHAIN (pos);
}
- bool changed = targetm.goacc.validate_dims (current_function_decl, dims);
+ bool changed = targetm.goacc.validate_dims (current_function_decl,
+ dims, fn_level);
- /* Default anything left undefaulted to 1. */
+ /* Default anything left to 1. */
for (ix = 0; ix != GOMP_DIM_MAX; ix++)
if (dims[ix] < 0)
{
- dims[ix] = (int)(dims[ix] < -1);
+ dims[ix] = 1;
changed = true;
}
-
+
if (changed)
{
/* Replace the attribute with new values. */
pos = NULL_TREE;
for (ix = GOMP_DIM_MAX; ix--;)
- pos = tree_cons (NULL_TREE,
+ pos = tree_cons (purpose[ix],
build_int_cst (integer_type_node, dims[ix]),
pos);
replace_oacc_fn_attrib (current_function_decl, pos);
@@ -14722,7 +14777,8 @@ execute_oacc_transform ()
hook. */
bool
-default_goacc_validate_dims (tree ARG_UNUSED (decl), int *ARG_UNUSED (dims))
+default_goacc_validate_dims (tree ARG_UNUSED (decl), int *ARG_UNUSED (dims),
+ int ARG_UNUSED (fn_level))
{
bool changed = false;
Index: gcc/omp-low.h
===================================================================
--- gcc/omp-low.h (revision 226860)
+++ gcc/omp-low.h (working copy)
@@ -30,6 +30,8 @@ extern bool make_gimple_omp_edges (basic
extern void omp_finish_file (void);
extern bool gimple_stmt_omp_data_i_init_p (gimple);
extern basic_block loop_get_oacc_kernels_region_entry (struct loop *);
+extern void replace_oacc_fn_attrib (tree, tree);
+extern tree build_oacc_routine_dims (tree);
extern tree get_oacc_fn_attrib (tree);
extern GTY(()) vec<tree, va_gc> *offload_funcs;
Index: gcc/target.def
===================================================================
--- gcc/target.def (revision 226860)
+++ gcc/target.def (working copy)
@@ -1651,7 +1651,7 @@ in anything that needs default to non-un
Defaults are represented as -1. Diagnostics should be issuedas \n\
ppropriate. Return true if changes have been made. You must override\n\
this hook to provide dimensions larger than 1.",
-bool, (tree, int []),
+bool, (tree, int [], int),
default_goacc_validate_dims)
DEFHOOK
Index: gcc/targhooks.h
===================================================================
--- gcc/targhooks.h (revision 226860)
+++ gcc/targhooks.h (working copy)
@@ -107,7 +107,7 @@ extern unsigned default_add_stmt_cost (v
extern void default_finish_cost (void *, unsigned *, unsigned *, unsigned *);
extern void default_destroy_cost_data (void *);
-extern bool default_goacc_validate_dims (tree, int []);
+extern bool default_goacc_validate_dims (tree, int [], int);
extern unsigned default_goacc_dim_limit (unsigned);
extern bool default_goacc_fork_join (gimple_stmt_iterator *, gimple,
const int [], bool);
Index: gcc/testsuite/c-c++-common/goacc/routine-4.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc/routine-4.c (revision 226860)
+++ gcc/testsuite/c-c++-common/goacc/routine-4.c (working copy)
@@ -1,87 +1,74 @@
/* Test invalid use of clauses with routine. */
/* { dg-do compile } */
-#pragma acc routine gang worker /* { dg-error "invalid combination" } */
+#pragma acc routine gang worker /* { dg-error "multiple loop axes" } */
void
f1 (void)
{
}
-#pragma acc routine worker gang /* { dg-error "invalid combination" } */
+#pragma acc routine worker gang /* { dg-error "multiple loop axes" } */
void
f1a (void)
{
}
-#pragma acc routine gang vector /* { dg-error "invalid combination" } */
+#pragma acc routine gang vector /* { dg-error "multiple loop axes" } */
void
f2 (void)
{
}
-#pragma acc routine vector gang /* { dg-error "invalid combination" } */
+#pragma acc routine vector gang /* { dg-error "multiple loop axes" } */
void
f2a (void)
{
}
-#pragma acc routine gang seq /* { dg-error "invalid combination" } */
+#pragma acc routine gang seq /* { dg-error "multiple loop axes" } */
void
f3 (void)
{
}
-#pragma acc routine seq gang /* { dg-error "invalid combination" } */
+#pragma acc routine seq gang /* { dg-error "multiple loop axes" } */
void
f3a (void)
{
}
-#pragma acc routine worker vector /* { dg-error "invalid combination" } */
+#pragma acc routine worker vector /* { dg-error "multiple loop axes" } */
void
f4 (void)
{
}
-#pragma acc routine vector worker /* { dg-error "invalid combination" } */
+#pragma acc routine vector worker /* { dg-error "multiple loop axes" } */
void
f4a (void)
{
}
-#pragma acc routine worker seq /* { dg-error "invalid combination" } */
+#pragma acc routine worker seq /* { dg-error "multiple loop axes" } */
void
f5 (void)
{
}
-#pragma acc routine seq worker /* { dg-error "invalid combination" } */
+#pragma acc routine seq worker /* { dg-error "multiple loop axes" } */
void
f5a (void)
{
}
-#pragma acc routine vector seq /* { dg-error "invalid combination" } */
+#pragma acc routine vector seq /* { dg-error "multiple loop axes" } */
void
f6 (void)
{
}
-#pragma acc routine seq vector /* { dg-error "invalid combination" } */
+#pragma acc routine seq vector /* { dg-error "multiple loop axes" } */
void
f6a (void)
{
}
-
-#pragma acc routine (g1) gang worker /* { dg-error "invalid combination" } */
-#pragma acc routine (g2) worker gang /* { dg-error "invalid combination" } */
-#pragma acc routine (g3) gang vector /* { dg-error "invalid combination" } */
-#pragma acc routine (g4) vector gang /* { dg-error "invalid combination" } */
-#pragma acc routine (g5) gang seq /* { dg-error "invalid combination" } */
-#pragma acc routine (g6) seq gang /* { dg-error "invalid combination" } */
-#pragma acc routine (g7) worker vector /* { dg-error "invalid combination" } */
-#pragma acc routine (g8) vector worker /* { dg-error "invalid combination" } */
-#pragma acc routine (g9) worker seq /* { dg-error "invalid combination" } */
-#pragma acc routine (g10) seq worker /* { dg-error "invalid combination" } */
-#pragma acc routine (g11) vector seq /* { dg-error "invalid combination" } */
-#pragma acc routine (g12) seq vector /* { dg-error "invalid combination" } */