2015-08-15 Nathan Sidwell * 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 *args) +set_oacc_fn_attrib (tree fn, tree clauses, vec *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 *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" } */