* [0/2] OpenACC routine support @ 2015-11-02 18:56 Nathan Sidwell 2015-11-02 19:21 ` [1/2] " Nathan Sidwell 2015-11-02 19:23 ` [2/2] " Nathan Sidwell 0 siblings, 2 replies; 23+ messages in thread From: Nathan Sidwell @ 2015-11-02 18:56 UTC (permalink / raw) To: Jakub Jelinek, GCC Patches The following pair of patches add support for routines. 01-trunk-routine-code-1102.patch: Compiler changes 02-trunk-routine-tests-1102.patch: Tests The changes are pretty straight forwards, as the execution model patch set laid the groundwork. Routines are decorated with the 'oacc routine' pragma, so they go through essentially the same processing as a parallels region, except of course they're already a device-side function. As with existing offloading the design is forward compatible with device_type support, which would operate in the same manner as an openacc offloaded region function WRT processing the attribute. nathan ^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [1/2] OpenACC routine support 2015-11-02 18:56 [0/2] OpenACC routine support Nathan Sidwell @ 2015-11-02 19:21 ` Nathan Sidwell 2015-11-03 15:35 ` Jakub Jelinek 2015-11-02 19:23 ` [2/2] " Nathan Sidwell 1 sibling, 1 reply; 23+ messages in thread From: Nathan Sidwell @ 2015-11-02 19:21 UTC (permalink / raw) To: Jakub Jelinek, GCC Patches [-- Attachment #1: Type: text/plain, Size: 837 bytes --] This is the core changes, an C & C++ FE parsing pieces. Parsing only deals with the gang, worker, vector & seq clauses. The nohost and bind clauses will be a later patch to port. The parsing is very similar to the omp declare simd parsing, in the way that it's hooked into the rest of the parser. conversion of the gang, worker, vecto & seq clauses to the internal 'oacc function' attribute format is done by a new routine, build_oacc_routine_dims, in omp-low.c. This sets the dimensions over which a routine might partition a loop to integer_zero_node, and sets TREE_PURPOSE of such dimensions to zero. while those dimensions over which it must not partition a loop have TREE_PURPOSE set to non-zero. the handling for validating this in the lower_oacc_device pass, and in the PTX backend is already there. ok? nathan [-- Attachment #2: 01-trunk-routine-code-1102.patch --] [-- Type: text/x-patch, Size: 27602 bytes --] 2015-11-02 Nathan Sidwell <nathan@codesourcery.com> * omp-low.h (replace_oacc_fn_attrib, build_oacc_routine_dims): Declare. * omp-low.c (build_oacc_routine_dims): New. 2015-11-02 Thomas Schwinge <thomas@codesourcery.com> Cesar Philippidis <cesar@codesourcery.com> James Norris <jnorris@codesourcery.com> Julian Brown <julian@codesourcery.com> Nathan Sidwell <nathan@codesourcery.com> c/ * c-parser.c (c_parser_declaration_or_fndef): Add OpenACC routine arg. Adjust all callers. (c_parser_declaration_or_fndef): Call c_finish_oacc_routine. (c_parser_pragma): Parse 'acc routine'. (OACC_ROUTINE_CLAUSE_MARK): Define. (c_parser_oacc_routine, (c_finish_oacc_routine): New. 2015-11-02 Thomas Schwinge <thomas@codesourcery.com> Cesar Philippidis <cesar@codesourcery.com> James Norris <jnorris@codesourcery.com> Julian Brown <julian@codesourcery.com> Nathan Sidwell <nathan@codesourcery.com> c-family/ * c-pragma.c (oacc_pragmas): Add "routine". * c-pragma.h (pragma_kind): Add PRAGMA_OACC_ROUTINE. 2015-11-02 Thomas Schwinge <thomas@codesourcery.com> Cesar Philippidis <cesar@codesourcery.com> James Norris <jnorris@codesourcery.com> Julian Brown <julian@codesourcery.com> Nathan Sidwell <nathan@codesourcery.com> cp/ * parser.h (struct cp_parser): Add oacc_routine field. * parser.c (cp_ensure_no_oacc_routine): New. (cp_parser_new): Initialize oacc_routine field. (cp_parser_linkage_specification): Call cp_ensure_no_oacc_routine. (cp_parser_namespace_definition, cp_parser_class_specifier_1): Likewise. (cp_parser_init_declarator): Call cp_finalize_oacc_routine. (cp_parser_function_definition, cp_parser_save_member_function_body): Likewise. (OACC_ROUTINE_CLAUSE_MASK): New. (cp_parser_finish_oacc_routine, cp_parser_oacc_routine, cp_finalize_oacc_routine): New. (cp_parser_pragma): Adjust omp_declare_simd checking. Call cp_ensure_no_oacc_routine. (cp_parser_pragma): Add OpenACC routine handling. Index: gcc/c/c-parser.c =================================================================== --- gcc/c/c-parser.c (revision 229667) +++ gcc/c/c-parser.c (working copy) @@ -1160,7 +1160,8 @@ enum c_parser_prec { static void c_parser_external_declaration (c_parser *); static void c_parser_asm_definition (c_parser *); static void c_parser_declaration_or_fndef (c_parser *, bool, bool, bool, - bool, bool, tree *, vec<c_token>); + bool, bool, tree *, vec<c_token>, + tree); static void c_parser_static_assert_declaration_no_semi (c_parser *); static void c_parser_static_assert_declaration (c_parser *); static void c_parser_declspecs (c_parser *, struct c_declspecs *, bool, bool, @@ -1249,6 +1250,7 @@ static bool c_parser_omp_target (c_parse static void c_parser_omp_end_declare_target (c_parser *); static void c_parser_omp_declare (c_parser *, enum pragma_context); static bool c_parser_omp_ordered (c_parser *, enum pragma_context); +static void c_parser_oacc_routine (c_parser *parser, enum pragma_context); /* These Objective-C parser functions are only ever called when compiling Objective-C. */ @@ -1428,12 +1430,13 @@ c_parser_external_declaration (c_parser only tell which after parsing the declaration specifiers, if any, and the first declarator. */ c_parser_declaration_or_fndef (parser, true, true, true, false, true, - NULL, vNULL); + NULL, vNULL, NULL_TREE); break; } } static void c_finish_omp_declare_simd (c_parser *, tree, tree, vec<c_token>); +static void c_finish_oacc_routine (c_parser *, tree, tree, bool, bool); /* Parse a declaration or function definition (C90 6.5, 6.7.1, C99 6.7, 6.9.1). If FNDEF_OK is true, a function definition is @@ -1511,7 +1514,8 @@ c_parser_declaration_or_fndef (c_parser bool static_assert_ok, bool empty_ok, bool nested, bool start_attr_ok, tree *objc_foreach_object_declaration, - vec<c_token> omp_declare_simd_clauses) + vec<c_token> omp_declare_simd_clauses, + tree oacc_routine_clauses) { struct c_declspecs *specs; tree prefix_attrs; @@ -1581,6 +1585,9 @@ c_parser_declaration_or_fndef (c_parser pedwarn (here, 0, "empty declaration"); } c_parser_consume_token (parser); + if (oacc_routine_clauses) + c_finish_oacc_routine (parser, NULL_TREE, + oacc_routine_clauses, false, false); return; } @@ -1697,6 +1704,9 @@ c_parser_declaration_or_fndef (c_parser || !vec_safe_is_empty (parser->cilk_simd_fn_tokens)) c_finish_omp_declare_simd (parser, NULL_TREE, NULL_TREE, omp_declare_simd_clauses); + if (oacc_routine_clauses) + c_finish_oacc_routine (parser, NULL_TREE, + oacc_routine_clauses, false, false); c_parser_skip_to_end_of_block_or_statement (parser); return; } @@ -1811,6 +1821,9 @@ c_parser_declaration_or_fndef (c_parser init = c_parser_initializer (parser); finish_init (); } + if (oacc_routine_clauses) + c_finish_oacc_routine (parser, d, oacc_routine_clauses, + false, false); if (d != error_mark_node) { maybe_warn_string_init (init_loc, TREE_TYPE (d), init); @@ -1854,6 +1867,9 @@ c_parser_declaration_or_fndef (c_parser if (parms) temp_pop_parm_decls (); } + if (oacc_routine_clauses) + c_finish_oacc_routine (parser, d, oacc_routine_clauses, + false, false); if (d) finish_decl (d, UNKNOWN_LOCATION, NULL_TREE, NULL_TREE, asm_name); @@ -1958,12 +1974,15 @@ c_parser_declaration_or_fndef (c_parser while (c_parser_next_token_is_not (parser, CPP_EOF) && c_parser_next_token_is_not (parser, CPP_OPEN_BRACE)) c_parser_declaration_or_fndef (parser, false, false, false, - true, false, NULL, vNULL); + true, false, NULL, vNULL, NULL_TREE); store_parm_decls (); if (omp_declare_simd_clauses.exists () || !vec_safe_is_empty (parser->cilk_simd_fn_tokens)) c_finish_omp_declare_simd (parser, current_function_decl, NULL_TREE, omp_declare_simd_clauses); + if (oacc_routine_clauses) + c_finish_oacc_routine (parser, current_function_decl, + oacc_routine_clauses, false, true); DECL_STRUCT_FUNCTION (current_function_decl)->function_start_locus = c_parser_peek_token (parser)->location; fnbody = c_parser_compound_statement (parser); @@ -4634,7 +4653,7 @@ c_parser_compound_statement_nostart (c_p last_label = false; mark_valid_location_for_stdc_pragma (false); c_parser_declaration_or_fndef (parser, true, true, true, true, - true, NULL, vNULL); + true, NULL, vNULL, NULL_TREE); if (last_stmt) pedwarn_c90 (loc, OPT_Wdeclaration_after_statement, "ISO C90 forbids mixed declarations and code"); @@ -4659,7 +4678,7 @@ c_parser_compound_statement_nostart (c_p last_label = false; mark_valid_location_for_stdc_pragma (false); c_parser_declaration_or_fndef (parser, true, true, true, true, - true, NULL, vNULL); + true, NULL, vNULL, NULL_TREE); /* Following the old parser, __extension__ does not disable this diagnostic. */ restore_extension_diagnostics (ext); @@ -4808,7 +4827,7 @@ c_parser_label (c_parser *parser) /*static_assert_ok*/ true, /*empty_ok*/ true, /*nested*/ true, /*start_attr_ok*/ true, NULL, - vNULL); + vNULL, NULL_TREE); } } } @@ -5580,7 +5599,7 @@ c_parser_for_statement (c_parser *parser else if (c_parser_next_tokens_start_declaration (parser)) { c_parser_declaration_or_fndef (parser, true, true, true, true, true, - &object_expression, vNULL); + &object_expression, vNULL, NULL_TREE); parser->objc_could_be_foreach_context = false; if (c_parser_next_token_is_keyword (parser, RID_IN)) @@ -5609,7 +5628,8 @@ c_parser_for_statement (c_parser *parser ext = disable_extension_diagnostics (); c_parser_consume_token (parser); c_parser_declaration_or_fndef (parser, true, true, true, true, - true, &object_expression, vNULL); + true, &object_expression, vNULL, + NULL_TREE); parser->objc_could_be_foreach_context = false; restore_extension_diagnostics (ext); @@ -8745,8 +8765,8 @@ c_parser_objc_methodprotolist (c_parser c_parser_consume_token (parser); } else - c_parser_declaration_or_fndef (parser, false, false, true, - false, true, NULL, vNULL); + c_parser_declaration_or_fndef (parser, false, false, true, false, + true, NULL, vNULL, NULL_TREE); break; } } @@ -9703,6 +9723,10 @@ c_parser_pragma (c_parser *parser, enum c_parser_oacc_enter_exit_data (parser, false); return false; + case PRAGMA_OACC_ROUTINE: + c_parser_oacc_routine (parser, context); + return false; + case PRAGMA_OACC_UPDATE: if (context != pragma_compound) { @@ -13265,6 +13289,111 @@ c_parser_oacc_kernels_parallel (location } /* OpenACC 2.0: + # pragma acc routine oacc-routine-clause[optseq] new-line + function-definition + + # pragma acc routine ( name ) oacc-routine-clause[optseq] new-line +*/ + +#define OACC_ROUTINE_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_GANG) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WORKER) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ) ) + +/* Parse an OpenACC routine directive. For named directives, we apply + immediately to the named function. For unnamed ones we then parse + a declaration or definition, which must be for a function. */ + +static void +c_parser_oacc_routine (c_parser *parser, enum pragma_context context) +{ + tree decl = NULL_TREE; + /* Create a dummy claue, to record location. */ + tree c_head = build_omp_clause (c_parser_peek_token (parser)->location, + OMP_CLAUSE_SEQ); + + if (context != pragma_external) + c_parser_error (parser, "%<#pragma acc routine%> not at file scope"); + + c_parser_consume_pragma (parser); + + /* Scan for optional '( name )'. */ + if (c_parser_peek_token (parser)->type == CPP_OPEN_PAREN) + { + c_parser_consume_token (parser); + + c_token *token = c_parser_peek_token (parser); + + if (token->type == CPP_NAME && token->id_kind == C_ID_ID) + { + decl = lookup_name (token->value); + if (!decl) + { + error_at (token->location, "%qE has not been declared", + token->value); + decl = error_mark_node; + } + c_parser_consume_token (parser); + } + else + c_parser_error (parser, "expected function name"); + + c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, 0); + } + + /* Build a chain of clauses. */ + parser->in_pragma = true; + tree clauses = c_parser_oacc_all_clauses + (parser, OACC_ROUTINE_CLAUSE_MASK, "#pragma acc routine"); + + /* Force clauses to be non-null, by attaching context to it. */ + clauses = tree_cons (c_head, clauses, NULL_TREE); + + if (decl) + c_finish_oacc_routine (parser, decl, clauses, true, false); + else + c_parser_declaration_or_fndef (parser, true, false, false, false, + true, NULL, vNULL, clauses); +} + +/* Finalize an OpenACC routine pragma, applying it to FNDECL. CLAUSES + are the parsed clauses. IS_DEFN is true if we're applying it to + the definition (so expect FNDEF to look somewhat defined. */ + +static void +c_finish_oacc_routine (c_parser *ARG_UNUSED (parser), + tree fndecl, tree clauses, bool named, bool is_defn) +{ + location_t loc = OMP_CLAUSE_LOCATION (TREE_PURPOSE (clauses)); + + if (!fndecl || TREE_CODE (fndecl) != FUNCTION_DECL) + { + if (fndecl != error_mark_node) + error_at (loc, "%<#pragma acc routine%> %s", + named ? "does not refer to a function" + : "not followed by function"); + return; + } + + if (get_oacc_fn_attrib (fndecl)) + error_at (loc, "%<#pragma acc routine%> already applied to %D", fndecl); + + if (TREE_USED (fndecl) || (!is_defn && DECL_SAVED_TREE (fndecl))) + error_at (loc, "%<#pragma acc routine%> must be applied before %s", + TREE_USED (fndecl) ? "use" : "definition"); + + /* Process for function attrib */ + tree dims = build_oacc_routine_dims (TREE_VALUE (clauses)); + replace_oacc_fn_attrib (fndecl, dims); + + /* Also attach as a declare. */ + DECL_ATTRIBUTES (fndecl) + = tree_cons (get_identifier ("omp declare target"), + clauses, DECL_ATTRIBUTES (fndecl)); +} + +/* OpenACC 2.0: # pragma acc update oacc-update-clause[optseq] new-line */ @@ -13929,7 +14058,7 @@ c_parser_omp_for_loop (location_t loc, c vec_safe_push (for_block, c_begin_compound_stmt (true)); this_pre_body = push_stmt_list (); c_parser_declaration_or_fndef (parser, true, true, true, true, true, - NULL, vNULL); + NULL, vNULL, NULL_TREE); if (this_pre_body) { this_pre_body = pop_stmt_list (this_pre_body); @@ -15506,12 +15635,12 @@ c_parser_omp_declare_simd (c_parser *par while (c_parser_next_token_is (parser, CPP_KEYWORD) && c_parser_peek_token (parser)->keyword == RID_EXTENSION); c_parser_declaration_or_fndef (parser, true, true, true, false, true, - NULL, clauses); + NULL, clauses, NULL_TREE); restore_extension_diagnostics (ext); } else c_parser_declaration_or_fndef (parser, true, true, true, false, true, - NULL, clauses); + NULL, clauses, NULL_TREE); break; case pragma_struct: case pragma_param: @@ -15531,7 +15660,7 @@ c_parser_omp_declare_simd (c_parser *par if (c_parser_next_tokens_start_declaration (parser)) { c_parser_declaration_or_fndef (parser, true, true, true, true, - true, NULL, clauses); + true, NULL, clauses, NULL_TREE); restore_extension_diagnostics (ext); break; } @@ -15540,7 +15669,7 @@ c_parser_omp_declare_simd (c_parser *par else if (c_parser_next_tokens_start_declaration (parser)) { c_parser_declaration_or_fndef (parser, true, true, true, true, true, - NULL, clauses); + NULL, clauses, NULL_TREE); break; } c_parser_error (parser, "%<#pragma omp declare simd%> must be followed by " Index: gcc/c-family/c-pragma.c =================================================================== --- gcc/c-family/c-pragma.c (revision 229667) +++ gcc/c-family/c-pragma.c (working copy) @@ -1211,6 +1211,7 @@ static const struct omp_pragma_def oacc_ { "kernels", PRAGMA_OACC_KERNELS }, { "loop", PRAGMA_OACC_LOOP }, { "parallel", PRAGMA_OACC_PARALLEL }, + { "routine", PRAGMA_OACC_ROUTINE }, { "update", PRAGMA_OACC_UPDATE }, { "wait", PRAGMA_OACC_WAIT } }; Index: gcc/c-family/c-pragma.h =================================================================== --- gcc/c-family/c-pragma.h (revision 229667) +++ gcc/c-family/c-pragma.h (working copy) @@ -34,6 +34,7 @@ enum pragma_kind { PRAGMA_OACC_KERNELS, PRAGMA_OACC_LOOP, PRAGMA_OACC_PARALLEL, + PRAGMA_OACC_ROUTINE, PRAGMA_OACC_UPDATE, PRAGMA_OACC_WAIT, Index: gcc/cp/parser.c =================================================================== --- gcc/cp/parser.c (revision 229667) +++ gcc/cp/parser.c (working copy) @@ -244,6 +244,8 @@ static bool cp_parser_omp_declare_reduct (tree, cp_parser *); static tree cp_parser_cilk_simd_vectorlength (cp_parser *, tree, bool); +static void cp_finalize_oacc_routine + (cp_parser *, tree, bool); /* Manifest constants. */ #define CP_LEXER_BUFFER_SIZE ((256 * 1024) / sizeof (cp_token)) @@ -1319,6 +1321,15 @@ cp_finalize_omp_declare_simd (cp_parser } } } + +/* Diagnose if #pragma omp routine isn't followed immediately + by function declaration or definition. */ + +static inline void +cp_ensure_no_oacc_routine (cp_parser *parser) +{ + cp_finalize_oacc_routine (parser, NULL_TREE, false); +} \f /* Decl-specifiers. */ @@ -3619,6 +3630,9 @@ cp_parser_new (void) parser->implicit_template_parms = 0; parser->implicit_template_scope = 0; + /* Active OpenACC routine clauses. */ + parser->oacc_routine = NULL; + /* Allow constrained-type-specifiers. */ parser->prevent_constrained_type_specifiers = 0; @@ -12535,6 +12549,7 @@ cp_parser_linkage_specification (cp_pars if (cp_lexer_next_token_is (parser->lexer, CPP_OPEN_BRACE)) { cp_ensure_no_omp_declare_simd (parser); + cp_ensure_no_oacc_routine (parser); /* Consume the `{' token. */ cp_lexer_consume_token (parser->lexer); @@ -17062,6 +17077,7 @@ cp_parser_namespace_definition (cp_parse int nested_definition_count = 0; cp_ensure_no_omp_declare_simd (parser); + cp_ensure_no_oacc_routine (parser); if (cp_lexer_next_token_is_keyword (parser->lexer, RID_INLINE)) { maybe_warn_cpp0x (CPP0X_INLINE_NAMESPACES); @@ -18085,6 +18101,7 @@ cp_parser_init_declarator (cp_parser* pa range_for_decl_p? SD_INITIALIZED : is_initialized, attributes, prefix_attributes, &pushed_scope); cp_finalize_omp_declare_simd (parser, decl); + cp_finalize_oacc_routine (parser, decl, false); /* Adjust location of decl if declarator->id_loc is more appropriate: set, and decl wasn't merged with another decl, in which case its location would be different from input_location, and more accurate. */ @@ -18198,6 +18215,7 @@ cp_parser_init_declarator (cp_parser* pa if (decl && TREE_CODE (decl) == FUNCTION_DECL) cp_parser_save_default_args (parser, decl); cp_finalize_omp_declare_simd (parser, decl); + cp_finalize_oacc_routine (parser, decl, false); } /* Finish processing the declaration. But, skip member @@ -20804,6 +20822,7 @@ cp_parser_class_specifier_1 (cp_parser* } cp_ensure_no_omp_declare_simd (parser); + cp_ensure_no_oacc_routine (parser); /* Issue an error message if type-definitions are forbidden here. */ cp_parser_check_type_definition (parser); @@ -22117,6 +22136,7 @@ cp_parser_member_declaration (cp_parser* } cp_finalize_omp_declare_simd (parser, decl); + cp_finalize_oacc_routine (parser, decl, false); /* Reset PREFIX_ATTRIBUTES. */ while (attributes && TREE_CHAIN (attributes) != first_attribute) @@ -24720,6 +24740,7 @@ cp_parser_function_definition_from_speci { cp_finalize_omp_declare_simd (parser, current_function_decl); parser->omp_declare_simd = NULL; + cp_finalize_oacc_routine (parser, current_function_decl, true); } if (!success_p) @@ -25402,6 +25423,7 @@ cp_parser_save_member_function_body (cp_ /* Create the FUNCTION_DECL. */ fn = grokmethod (decl_specifiers, declarator, attributes); cp_finalize_omp_declare_simd (parser, fn); + cp_finalize_oacc_routine (parser, fn, true); /* If something went badly wrong, bail out now. */ if (fn == error_mark_node) { @@ -35453,6 +35475,147 @@ cp_parser_omp_taskloop (cp_parser *parse return ret; } + +/* OpenACC 2.0: + # pragma acc routine oacc-routine-clause[optseq] new-line + function-definition + + # pragma acc routine ( name ) oacc-routine-clause[optseq] new-line +*/ + +#define OACC_ROUTINE_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_GANG) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WORKER) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ)) + +/* Finalize #pragma acc routine clauses after direct declarator has + been parsed, and put that into "omp declare target" attribute. */ + +static void +cp_parser_finish_oacc_routine (cp_parser *ARG_UNUSED (parser), tree fndecl, + tree clauses, bool named, bool is_defn) +{ + location_t loc = OMP_CLAUSE_LOCATION (TREE_PURPOSE (clauses)); + + if (named && fndecl && is_overloaded_fn (fndecl) + && (TREE_CODE (fndecl) != FUNCTION_DECL + || DECL_FUNCTION_TEMPLATE_P (fndecl))) + { + error_at (loc, "%<#pragma acc routine%> names a set of overloads"); + return; + } + + if (!fndecl || TREE_CODE (fndecl) != FUNCTION_DECL) + { + error_at (loc, "%<#pragma acc routine%> %s", + named ? "does not refer to a function" + : "not followed by function"); + return; + } + + /* Perhaps we should use the same rule as declarations in different + namespaces? */ + if (named && !DECL_NAMESPACE_SCOPE_P (fndecl)) + { + error_at (loc, "%<#pragma acc routine%> does not refer to a" + " namespace scope function"); + return; + } + + if (get_oacc_fn_attrib (fndecl)) + error_at (loc, "%<#pragma acc routine%> already applied to %D", fndecl); + + if (TREE_USED (fndecl) || (!is_defn && DECL_SAVED_TREE (fndecl))) + error_at (OMP_CLAUSE_LOCATION (TREE_PURPOSE (clauses)), + "%<#pragma acc routine%> must be applied before %s", + TREE_USED (fndecl) ? "use" : "definition"); + + /* Process for function attrib */ + tree dims = build_oacc_routine_dims (TREE_VALUE (clauses)); + replace_oacc_fn_attrib (fndecl, dims); + + /* Also attach as a declare. */ + DECL_ATTRIBUTES (fndecl) + = tree_cons (get_identifier ("omp declare target"), + clauses, DECL_ATTRIBUTES (fndecl)); +} + +/* Parse the OpenACC routine pragma. This has an optional '( name )' + component, which must resolve to a declared namespace-scope + function. The clauses are either processed directly (for a named + function), or defered until the immediatley following declaration + is parsed. */ + +static void +cp_parser_oacc_routine (cp_parser *parser, cp_token *pragma_tok, + enum pragma_context context) +{ + tree decl = NULL_TREE; + /* Create a dummy claue, to record location. */ + tree c_head = build_omp_clause (pragma_tok->location, OMP_CLAUSE_SEQ); + + if (context != pragma_external) + cp_parser_error (parser, "%<#pragma acc routine%> not at file scope"); + + /* Look for optional '( name )'. */ + if (cp_lexer_next_token_is (parser->lexer,CPP_OPEN_PAREN)) + { + cp_lexer_consume_token (parser->lexer); + cp_token *token = cp_lexer_peek_token (parser->lexer); + + /* We parse the name as an id-expression. If it resolves to + anything other than a non-overloaded function at namespace + scope, it's an error. */ + tree id = cp_parser_id_expression (parser, + /*template_keyword_p=*/false, + /*check_dependency_p=*/false, + /*template_p=*/NULL, + /*declarator_p=*/false, + /*optional_p=*/false); + decl = cp_parser_lookup_name_simple (parser, id, token->location); + if (id != error_mark_node && decl == error_mark_node) + cp_parser_name_lookup_error (parser, id, decl, NLE_NULL, + token->location); + + if (decl == error_mark_node + || !cp_parser_require (parser, CPP_CLOSE_PAREN, RT_CLOSE_PAREN)) + { + cp_parser_skip_to_pragma_eol (parser, pragma_tok); + return; + } + } + + /* Build a chain of clauses. */ + parser->lexer->in_pragma = true; + tree clauses = NULL_TREE; + clauses = cp_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK, + "#pragma acc routine", + cp_lexer_peek_token (parser->lexer)); + + /* Force clauses to be non-null, by attaching context to it. */ + clauses = tree_cons (c_head, clauses, NULL_TREE); + + if (decl) + cp_parser_finish_oacc_routine (parser, decl, clauses, true, false); + else + parser->oacc_routine = clauses; +} + +/* Apply any saved OpenACC routine clauses to a just-parsed + declaration. */ + +static void +cp_finalize_oacc_routine (cp_parser *parser, tree fndecl, bool is_defn) +{ + if (parser->oacc_routine) + { + cp_parser_finish_oacc_routine (parser, fndecl, parser->oacc_routine, + false, is_defn); + parser->oacc_routine = NULL_TREE; + } +} + /* Main entry point to OpenMP statement pragmas. */ static void @@ -35929,8 +36092,9 @@ cp_parser_pragma (cp_parser *parser, enu parser->lexer->in_pragma = true; id = pragma_tok->pragma_kind; - if (id != PRAGMA_OMP_DECLARE_REDUCTION) + if (id != PRAGMA_OMP_DECLARE_REDUCTION && id != PRAGMA_OACC_ROUTINE) cp_ensure_no_omp_declare_simd (parser); + cp_ensure_no_oacc_routine (parser); switch (id) { case PRAGMA_GCC_PCH_PREPROCESS: @@ -36040,6 +36204,10 @@ cp_parser_pragma (cp_parser *parser, enu cp_parser_omp_declare (parser, pragma_tok, context); return false; + case PRAGMA_OACC_ROUTINE: + cp_parser_oacc_routine (parser, pragma_tok, context); + return false; + case PRAGMA_OACC_CACHE: case PRAGMA_OACC_DATA: case PRAGMA_OACC_ENTER_DATA: Index: gcc/cp/parser.h =================================================================== --- gcc/cp/parser.h (revision 229667) +++ gcc/cp/parser.h (working copy) @@ -371,6 +371,9 @@ struct GTY(()) cp_parser { necessary. */ cp_omp_declare_simd_data * GTY((skip)) cilk_simd_fn_info; + /* OpenACC routine clauses for subsequent decl/defn. */ + tree oacc_routine; + /* Nonzero if parsing a parameter list where 'auto' should trigger an implicit template parameter. */ bool auto_is_implicit_function_template_parm_p; Index: gcc/omp-low.c =================================================================== --- gcc/omp-low.c (revision 229667) +++ gcc/omp-low.c (working copy) @@ -12083,6 +12083,50 @@ set_oacc_fn_attrib (tree fn, tree clause } } +/* 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. non-zero indicates + yes, zero indicates no. By construction once a non-zero has been + reached, further inner dimensions must also be non-zero. We set + TREE_VALUE to zero for the dimensions that may be partitioned and + 1 for the other ones -- if a loop is (erroneously) spawned at + an outer level, we don't want to try and partition it. */ + +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; + } + + /* Default to SEQ. */ + 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), + build_int_cst (integer_type_node, ix < level), dims); + + return dims; +} + /* Retrieve the oacc function attrib and return it. Non-oacc functions will return NULL. */ Index: gcc/omp-low.h =================================================================== --- gcc/omp-low.h (revision 229667) +++ gcc/omp-low.h (working copy) @@ -30,6 +30,8 @@ extern tree omp_reduction_init (tree, tr 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 tree build_oacc_routine_dims (tree); extern tree get_oacc_fn_attrib (tree); extern GTY(()) vec<tree, va_gc> *offload_funcs; ^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [1/2] OpenACC routine support 2015-11-02 19:21 ` [1/2] " Nathan Sidwell @ 2015-11-03 15:35 ` Jakub Jelinek 2015-11-03 15:55 ` Nathan Sidwell 2015-11-10 0:32 ` Nathan Sidwell 0 siblings, 2 replies; 23+ messages in thread From: Jakub Jelinek @ 2015-11-03 15:35 UTC (permalink / raw) To: Nathan Sidwell; +Cc: GCC Patches On Mon, Nov 02, 2015 at 02:21:43PM -0500, Nathan Sidwell wrote: > --- gcc/c/c-parser.c (revision 229667) > +++ gcc/c/c-parser.c (working copy) > @@ -1160,7 +1160,8 @@ enum c_parser_prec { > static void c_parser_external_declaration (c_parser *); > static void c_parser_asm_definition (c_parser *); > static void c_parser_declaration_or_fndef (c_parser *, bool, bool, bool, > - bool, bool, tree *, vec<c_token>); > + bool, bool, tree *, vec<c_token>, > + tree); Wonder if this shouldn't be tree = NULL_TREE, then you'd avoid most of the c_parser_declaration_or_fndef caller changes. Otherwise, LGTM. Jakub ^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [1/2] OpenACC routine support 2015-11-03 15:35 ` Jakub Jelinek @ 2015-11-03 15:55 ` Nathan Sidwell 2015-11-10 0:32 ` Nathan Sidwell 1 sibling, 0 replies; 23+ messages in thread From: Nathan Sidwell @ 2015-11-03 15:55 UTC (permalink / raw) To: Jakub Jelinek; +Cc: GCC Patches On 11/03/15 10:35, Jakub Jelinek wrote: > On Mon, Nov 02, 2015 at 02:21:43PM -0500, Nathan Sidwell wrote: >> --- gcc/c/c-parser.c (revision 229667) >> +++ gcc/c/c-parser.c (working copy) >> @@ -1160,7 +1160,8 @@ enum c_parser_prec { >> static void c_parser_external_declaration (c_parser *); >> static void c_parser_asm_definition (c_parser *); >> static void c_parser_declaration_or_fndef (c_parser *, bool, bool, bool, >> - bool, bool, tree *, vec<c_token>); >> + bool, bool, tree *, vec<c_token>, >> + tree); > > Wonder if this shouldn't be tree = NULL_TREE, then you'd avoid most of the > c_parser_declaration_or_fndef caller changes. yeah, I guess that'd be less invasive. I'm fine with it. ^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [1/2] OpenACC routine support 2015-11-03 15:35 ` Jakub Jelinek 2015-11-03 15:55 ` Nathan Sidwell @ 2015-11-10 0:32 ` Nathan Sidwell 2015-11-10 0:48 ` Nathan Sidwell 2015-11-10 5:28 ` Cesar Philippidis 1 sibling, 2 replies; 23+ messages in thread From: Nathan Sidwell @ 2015-11-10 0:32 UTC (permalink / raw) To: Jakub Jelinek; +Cc: GCC Patches, Cesar Philippidis [-- Attachment #1: Type: text/plain, Size: 914 bytes --] On 11/03/15 10:35, Jakub Jelinek wrote: > On Mon, Nov 02, 2015 at 02:21:43PM -0500, Nathan Sidwell wrote: >> --- gcc/c/c-parser.c (revision 229667) >> +++ gcc/c/c-parser.c (working copy) >> @@ -1160,7 +1160,8 @@ enum c_parser_prec { >> static void c_parser_external_declaration (c_parser *); >> static void c_parser_asm_definition (c_parser *); >> static void c_parser_declaration_or_fndef (c_parser *, bool, bool, bool, >> - bool, bool, tree *, vec<c_token>); >> + bool, bool, tree *, vec<c_token>, >> + tree); > > Wonder if this shouldn't be tree = NULL_TREE, then you'd avoid most of the > c_parser_declaration_or_fndef caller changes. > > Otherwise, LGTM. This is the patch I've just committed. It includes c parser adjustments to detect the case of two function decls with a single type specifier. Cesar will be applying a patch for the C++ parser for the same case. nathan [-- Attachment #2: trunk-routine-1109.patch --] [-- Type: text/x-patch, Size: 23763 bytes --] 2015-11-09 Nathan Sidwell <nathan@codesourcery.com> * omp-low.h (replace_oacc_fn_attrib, build_oacc_routine_dims): Declare. * omp-low.c (build_oacc_routine_dims): New. 2015-11-09 Thomas Schwinge <thomas@codesourcery.com> Cesar Philippidis <cesar@codesourcery.com> James Norris <jnorris@codesourcery.com> Julian Brown <julian@codesourcery.com> Nathan Sidwell <nathan@codesourcery.com> c/ * c-parser.c (c_parser_declaration_or_fndef): Add OpenACC routine arg. (c_parser_declaration_or_fndef): Call c_finish_oacc_routine. (c_parser_pragma): Parse 'acc routine'. (OACC_ROUTINE_CLAUSE_MARK): Define. (c_parser_oacc_routine, (c_finish_oacc_routine): New. 2015-11-09 Thomas Schwinge <thomas@codesourcery.com> Cesar Philippidis <cesar@codesourcery.com> James Norris <jnorris@codesourcery.com> Julian Brown <julian@codesourcery.com> Nathan Sidwell <nathan@codesourcery.com> c-family/ * c-pragma.c (oacc_pragmas): Add "routine". * c-pragma.h (pragma_kind): Add PRAGMA_OACC_ROUTINE. 2015-11-09 Thomas Schwinge <thomas@codesourcery.com> Cesar Philippidis <cesar@codesourcery.com> James Norris <jnorris@codesourcery.com> Julian Brown <julian@codesourcery.com> Nathan Sidwell <nathan@codesourcery.com> cp/ * parser.h (struct cp_parser): Add oacc_routine field. * parser.c (cp_ensure_no_oacc_routine): New. (cp_parser_new): Initialize oacc_routine field. (cp_parser_linkage_specification): Call cp_ensure_no_oacc_routine. (cp_parser_namespace_definition, cp_parser_class_specifier_1): Likewise. (cp_parser_init_declarator): Call cp_finalize_oacc_routine. (cp_parser_function_definition, cp_parser_save_member_function_body): Likewise. (OACC_ROUTINE_CLAUSE_MASK): New. (cp_parser_finish_oacc_routine, cp_parser_oacc_routine, cp_finalize_oacc_routine): New. (cp_parser_pragma): Adjust omp_declare_simd checking. Call cp_ensure_no_oacc_routine. (cp_parser_pragma): Add OpenACC routine handling. Index: gcc/omp-low.c =================================================================== --- gcc/omp-low.c (revision 230040) +++ gcc/omp-low.c (working copy) @@ -12361,6 +12361,50 @@ set_oacc_fn_attrib (tree fn, tree clause } } +/* 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. non-zero indicates + yes, zero indicates no. By construction once a non-zero has been + reached, further inner dimensions must also be non-zero. We set + TREE_VALUE to zero for the dimensions that may be partitioned and + 1 for the other ones -- if a loop is (erroneously) spawned at + an outer level, we don't want to try and partition it. */ + +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; + } + + /* Default to SEQ. */ + 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), + build_int_cst (integer_type_node, ix < level), dims); + + return dims; +} + /* Retrieve the oacc function attrib and return it. Non-oacc functions will return NULL. */ Index: gcc/omp-low.h =================================================================== --- gcc/omp-low.h (revision 230040) +++ gcc/omp-low.h (working copy) @@ -30,6 +30,8 @@ extern tree omp_reduction_init (tree, tr 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 tree build_oacc_routine_dims (tree); extern tree get_oacc_fn_attrib (tree); extern int get_oacc_ifn_dim_arg (const gimple *); extern int get_oacc_fn_dim_size (tree, int); Index: gcc/c-family/c-pragma.c =================================================================== --- gcc/c-family/c-pragma.c (revision 230040) +++ gcc/c-family/c-pragma.c (working copy) @@ -1257,6 +1257,7 @@ static const struct omp_pragma_def oacc_ { "kernels", PRAGMA_OACC_KERNELS }, { "loop", PRAGMA_OACC_LOOP }, { "parallel", PRAGMA_OACC_PARALLEL }, + { "routine", PRAGMA_OACC_ROUTINE }, { "update", PRAGMA_OACC_UPDATE }, { "wait", PRAGMA_OACC_WAIT } }; Index: gcc/c-family/c-pragma.h =================================================================== --- gcc/c-family/c-pragma.h (revision 230040) +++ gcc/c-family/c-pragma.h (working copy) @@ -35,6 +35,7 @@ enum pragma_kind { PRAGMA_OACC_KERNELS, PRAGMA_OACC_LOOP, PRAGMA_OACC_PARALLEL, + PRAGMA_OACC_ROUTINE, PRAGMA_OACC_UPDATE, PRAGMA_OACC_WAIT, Index: gcc/c/c-parser.c =================================================================== --- gcc/c/c-parser.c (revision 230040) +++ gcc/c/c-parser.c (working copy) @@ -1162,7 +1162,8 @@ enum c_parser_prec { static void c_parser_external_declaration (c_parser *); static void c_parser_asm_definition (c_parser *); static void c_parser_declaration_or_fndef (c_parser *, bool, bool, bool, - bool, bool, tree *, vec<c_token>); + bool, bool, tree *, vec<c_token>, + tree = NULL_TREE); static void c_parser_static_assert_declaration_no_semi (c_parser *); static void c_parser_static_assert_declaration (c_parser *); static void c_parser_declspecs (c_parser *, struct c_declspecs *, bool, bool, @@ -1251,6 +1252,7 @@ static bool c_parser_omp_target (c_parse static void c_parser_omp_end_declare_target (c_parser *); static void c_parser_omp_declare (c_parser *, enum pragma_context); static bool c_parser_omp_ordered (c_parser *, enum pragma_context); +static void c_parser_oacc_routine (c_parser *parser, enum pragma_context); /* These Objective-C parser functions are only ever called when compiling Objective-C. */ @@ -1436,6 +1438,7 @@ c_parser_external_declaration (c_parser } static void c_finish_omp_declare_simd (c_parser *, tree, tree, vec<c_token>); +static void c_finish_oacc_routine (c_parser *, tree, tree, bool, bool, bool); /* Parse a declaration or function definition (C90 6.5, 6.7.1, C99 6.7, 6.9.1). If FNDEF_OK is true, a function definition is @@ -1513,7 +1516,8 @@ c_parser_declaration_or_fndef (c_parser bool static_assert_ok, bool empty_ok, bool nested, bool start_attr_ok, tree *objc_foreach_object_declaration, - vec<c_token> omp_declare_simd_clauses) + vec<c_token> omp_declare_simd_clauses, + tree oacc_routine_clauses) { struct c_declspecs *specs; tree prefix_attrs; @@ -1583,6 +1587,9 @@ c_parser_declaration_or_fndef (c_parser pedwarn (here, 0, "empty declaration"); } c_parser_consume_token (parser); + if (oacc_routine_clauses) + c_finish_oacc_routine (parser, NULL_TREE, + oacc_routine_clauses, false, true, false); return; } @@ -1680,7 +1687,7 @@ c_parser_declaration_or_fndef (c_parser prefix_attrs = specs->attrs; all_prefix_attrs = prefix_attrs; specs->attrs = NULL_TREE; - while (true) + for (bool first = true;; first = false) { struct c_declarator *declarator; bool dummy = false; @@ -1699,6 +1706,10 @@ c_parser_declaration_or_fndef (c_parser || !vec_safe_is_empty (parser->cilk_simd_fn_tokens)) c_finish_omp_declare_simd (parser, NULL_TREE, NULL_TREE, omp_declare_simd_clauses); + if (oacc_routine_clauses) + c_finish_oacc_routine (parser, NULL_TREE, + oacc_routine_clauses, + false, first, false); c_parser_skip_to_end_of_block_or_statement (parser); return; } @@ -1813,6 +1824,9 @@ c_parser_declaration_or_fndef (c_parser init = c_parser_initializer (parser); finish_init (); } + if (oacc_routine_clauses) + c_finish_oacc_routine (parser, d, oacc_routine_clauses, + false, first, false); if (d != error_mark_node) { maybe_warn_string_init (init_loc, TREE_TYPE (d), init); @@ -1856,6 +1870,9 @@ c_parser_declaration_or_fndef (c_parser if (parms) temp_pop_parm_decls (); } + if (oacc_routine_clauses) + c_finish_oacc_routine (parser, d, oacc_routine_clauses, + false, first, false); if (d) finish_decl (d, UNKNOWN_LOCATION, NULL_TREE, NULL_TREE, asm_name); @@ -1966,6 +1983,9 @@ c_parser_declaration_or_fndef (c_parser || !vec_safe_is_empty (parser->cilk_simd_fn_tokens)) c_finish_omp_declare_simd (parser, current_function_decl, NULL_TREE, omp_declare_simd_clauses); + if (oacc_routine_clauses) + c_finish_oacc_routine (parser, current_function_decl, + oacc_routine_clauses, false, first, true); DECL_STRUCT_FUNCTION (current_function_decl)->function_start_locus = c_parser_peek_token (parser)->location; fnbody = c_parser_compound_statement (parser); @@ -9706,6 +9726,10 @@ c_parser_pragma (c_parser *parser, enum c_parser_oacc_enter_exit_data (parser, false); return false; + case PRAGMA_OACC_ROUTINE: + c_parser_oacc_routine (parser, context); + return false; + case PRAGMA_OACC_UPDATE: if (context != pragma_compound) { @@ -13400,6 +13424,117 @@ c_parser_oacc_kernels_parallel (location } /* OpenACC 2.0: + # pragma acc routine oacc-routine-clause[optseq] new-line + function-definition + + # pragma acc routine ( name ) oacc-routine-clause[optseq] new-line +*/ + +#define OACC_ROUTINE_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_GANG) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WORKER) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ) ) + +/* Parse an OpenACC routine directive. For named directives, we apply + immediately to the named function. For unnamed ones we then parse + a declaration or definition, which must be for a function. */ + +static void +c_parser_oacc_routine (c_parser *parser, enum pragma_context context) +{ + tree decl = NULL_TREE; + /* Create a dummy claue, to record location. */ + tree c_head = build_omp_clause (c_parser_peek_token (parser)->location, + OMP_CLAUSE_SEQ); + + if (context != pragma_external) + c_parser_error (parser, "%<#pragma acc routine%> not at file scope"); + + c_parser_consume_pragma (parser); + + /* Scan for optional '( name )'. */ + if (c_parser_peek_token (parser)->type == CPP_OPEN_PAREN) + { + c_parser_consume_token (parser); + + c_token *token = c_parser_peek_token (parser); + + if (token->type == CPP_NAME && (token->id_kind == C_ID_ID + || token->id_kind == C_ID_TYPENAME)) + { + decl = lookup_name (token->value); + if (!decl) + { + error_at (token->location, "%qE has not been declared", + token->value); + decl = error_mark_node; + } + } + else + c_parser_error (parser, "expected function name"); + + if (token->type != CPP_CLOSE_PAREN) + c_parser_consume_token (parser); + + c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, 0); + } + + /* Build a chain of clauses. */ + parser->in_pragma = true; + tree clauses = c_parser_oacc_all_clauses + (parser, OACC_ROUTINE_CLAUSE_MASK, "#pragma acc routine"); + + /* Force clauses to be non-null, by attaching context to it. */ + clauses = tree_cons (c_head, clauses, NULL_TREE); + + if (decl) + c_finish_oacc_routine (parser, decl, clauses, true, true, false); + else if (c_parser_peek_token (parser)->type == CPP_PRAGMA) + /* This will emit an error. */ + c_finish_oacc_routine (parser, NULL_TREE, clauses, false, true, false); + else + c_parser_declaration_or_fndef (parser, true, false, false, false, + true, NULL, vNULL, clauses); +} + +/* Finalize an OpenACC routine pragma, applying it to FNDECL. CLAUSES + are the parsed clauses. IS_DEFN is true if we're applying it to + the definition (so expect FNDEF to look somewhat defined. */ + +static void +c_finish_oacc_routine (c_parser *ARG_UNUSED (parser), tree fndecl, + tree clauses, bool named, bool first, bool is_defn) +{ + location_t loc = OMP_CLAUSE_LOCATION (TREE_PURPOSE (clauses)); + + if (!fndecl || TREE_CODE (fndecl) != FUNCTION_DECL || !first) + { + if (fndecl != error_mark_node) + error_at (loc, "%<#pragma acc routine%> %s", + named ? "does not refer to a function" + : "not followed by single function"); + return; + } + + if (get_oacc_fn_attrib (fndecl)) + error_at (loc, "%<#pragma acc routine%> already applied to %D", fndecl); + + if (TREE_USED (fndecl) || (!is_defn && DECL_SAVED_TREE (fndecl))) + error_at (loc, "%<#pragma acc routine%> must be applied before %s", + TREE_USED (fndecl) ? "use" : "definition"); + + /* Process for function attrib */ + tree dims = build_oacc_routine_dims (TREE_VALUE (clauses)); + replace_oacc_fn_attrib (fndecl, dims); + + /* Also attach as a declare. */ + DECL_ATTRIBUTES (fndecl) + = tree_cons (get_identifier ("omp declare target"), + clauses, DECL_ATTRIBUTES (fndecl)); +} + +/* OpenACC 2.0: # pragma acc update oacc-update-clause[optseq] new-line */ Index: gcc/cp/parser.c =================================================================== --- gcc/cp/parser.c (revision 230040) +++ gcc/cp/parser.c (working copy) @@ -245,6 +245,8 @@ static bool cp_parser_omp_declare_reduct (tree, cp_parser *); static tree cp_parser_cilk_simd_vectorlength (cp_parser *, tree, bool); +static void cp_finalize_oacc_routine + (cp_parser *, tree, bool); /* Manifest constants. */ #define CP_LEXER_BUFFER_SIZE ((256 * 1024) / sizeof (cp_token)) @@ -1320,6 +1322,15 @@ cp_finalize_omp_declare_simd (cp_parser } } } + +/* Diagnose if #pragma omp routine isn't followed immediately + by function declaration or definition. */ + +static inline void +cp_ensure_no_oacc_routine (cp_parser *parser) +{ + cp_finalize_oacc_routine (parser, NULL_TREE, false); +} \f /* Decl-specifiers. */ @@ -3620,6 +3631,9 @@ cp_parser_new (void) parser->implicit_template_parms = 0; parser->implicit_template_scope = 0; + /* Active OpenACC routine clauses. */ + parser->oacc_routine = NULL; + /* Allow constrained-type-specifiers. */ parser->prevent_constrained_type_specifiers = 0; @@ -12541,6 +12555,7 @@ cp_parser_linkage_specification (cp_pars if (cp_lexer_next_token_is (parser->lexer, CPP_OPEN_BRACE)) { cp_ensure_no_omp_declare_simd (parser); + cp_ensure_no_oacc_routine (parser); /* Consume the `{' token. */ cp_lexer_consume_token (parser->lexer); @@ -17058,6 +17073,7 @@ cp_parser_namespace_definition (cp_parse int nested_definition_count = 0; cp_ensure_no_omp_declare_simd (parser); + cp_ensure_no_oacc_routine (parser); if (cp_lexer_next_token_is_keyword (parser->lexer, RID_INLINE)) { maybe_warn_cpp0x (CPP0X_INLINE_NAMESPACES); @@ -18081,6 +18097,7 @@ cp_parser_init_declarator (cp_parser* pa range_for_decl_p? SD_INITIALIZED : is_initialized, attributes, prefix_attributes, &pushed_scope); cp_finalize_omp_declare_simd (parser, decl); + cp_finalize_oacc_routine (parser, decl, false); /* Adjust location of decl if declarator->id_loc is more appropriate: set, and decl wasn't merged with another decl, in which case its location would be different from input_location, and more accurate. */ @@ -18194,6 +18211,7 @@ cp_parser_init_declarator (cp_parser* pa if (decl && TREE_CODE (decl) == FUNCTION_DECL) cp_parser_save_default_args (parser, decl); cp_finalize_omp_declare_simd (parser, decl); + cp_finalize_oacc_routine (parser, decl, false); } /* Finish processing the declaration. But, skip member @@ -20800,6 +20818,7 @@ cp_parser_class_specifier_1 (cp_parser* } cp_ensure_no_omp_declare_simd (parser); + cp_ensure_no_oacc_routine (parser); /* Issue an error message if type-definitions are forbidden here. */ cp_parser_check_type_definition (parser); @@ -22113,6 +22132,7 @@ cp_parser_member_declaration (cp_parser* } cp_finalize_omp_declare_simd (parser, decl); + cp_finalize_oacc_routine (parser, decl, false); /* Reset PREFIX_ATTRIBUTES. */ while (attributes && TREE_CHAIN (attributes) != first_attribute) @@ -24716,6 +24736,7 @@ cp_parser_function_definition_from_speci { cp_finalize_omp_declare_simd (parser, current_function_decl); parser->omp_declare_simd = NULL; + cp_finalize_oacc_routine (parser, current_function_decl, true); } if (!success_p) @@ -25398,6 +25419,7 @@ cp_parser_save_member_function_body (cp_ /* Create the FUNCTION_DECL. */ fn = grokmethod (decl_specifiers, declarator, attributes); cp_finalize_omp_declare_simd (parser, fn); + cp_finalize_oacc_routine (parser, fn, true); /* If something went badly wrong, bail out now. */ if (fn == error_mark_node) { @@ -35584,6 +35606,147 @@ cp_parser_omp_taskloop (cp_parser *parse return ret; } + +/* OpenACC 2.0: + # pragma acc routine oacc-routine-clause[optseq] new-line + function-definition + + # pragma acc routine ( name ) oacc-routine-clause[optseq] new-line +*/ + +#define OACC_ROUTINE_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_GANG) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WORKER) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ)) + +/* Finalize #pragma acc routine clauses after direct declarator has + been parsed, and put that into "omp declare target" attribute. */ + +static void +cp_parser_finish_oacc_routine (cp_parser *ARG_UNUSED (parser), tree fndecl, + tree clauses, bool named, bool is_defn) +{ + location_t loc = OMP_CLAUSE_LOCATION (TREE_PURPOSE (clauses)); + + if (named && fndecl && is_overloaded_fn (fndecl) + && (TREE_CODE (fndecl) != FUNCTION_DECL + || DECL_FUNCTION_TEMPLATE_P (fndecl))) + { + error_at (loc, "%<#pragma acc routine%> names a set of overloads"); + return; + } + + if (!fndecl || TREE_CODE (fndecl) != FUNCTION_DECL) + { + error_at (loc, "%<#pragma acc routine%> %s", + named ? "does not refer to a function" + : "not followed by single function"); + return; + } + + /* Perhaps we should use the same rule as declarations in different + namespaces? */ + if (named && !DECL_NAMESPACE_SCOPE_P (fndecl)) + { + error_at (loc, "%<#pragma acc routine%> does not refer to a" + " namespace scope function"); + return; + } + + if (get_oacc_fn_attrib (fndecl)) + error_at (loc, "%<#pragma acc routine%> already applied to %D", fndecl); + + if (TREE_USED (fndecl) || (!is_defn && DECL_SAVED_TREE (fndecl))) + error_at (OMP_CLAUSE_LOCATION (TREE_PURPOSE (clauses)), + "%<#pragma acc routine%> must be applied before %s", + TREE_USED (fndecl) ? "use" : "definition"); + + /* Process for function attrib */ + tree dims = build_oacc_routine_dims (TREE_VALUE (clauses)); + replace_oacc_fn_attrib (fndecl, dims); + + /* Also attach as a declare. */ + DECL_ATTRIBUTES (fndecl) + = tree_cons (get_identifier ("omp declare target"), + clauses, DECL_ATTRIBUTES (fndecl)); +} + +/* Parse the OpenACC routine pragma. This has an optional '( name )' + component, which must resolve to a declared namespace-scope + function. The clauses are either processed directly (for a named + function), or defered until the immediatley following declaration + is parsed. */ + +static void +cp_parser_oacc_routine (cp_parser *parser, cp_token *pragma_tok, + enum pragma_context context) +{ + tree decl = NULL_TREE; + /* Create a dummy claue, to record location. */ + tree c_head = build_omp_clause (pragma_tok->location, OMP_CLAUSE_SEQ); + + if (context != pragma_external) + cp_parser_error (parser, "%<#pragma acc routine%> not at file scope"); + + /* Look for optional '( name )'. */ + if (cp_lexer_next_token_is (parser->lexer,CPP_OPEN_PAREN)) + { + cp_lexer_consume_token (parser->lexer); + cp_token *token = cp_lexer_peek_token (parser->lexer); + + /* We parse the name as an id-expression. If it resolves to + anything other than a non-overloaded function at namespace + scope, it's an error. */ + tree id = cp_parser_id_expression (parser, + /*template_keyword_p=*/false, + /*check_dependency_p=*/false, + /*template_p=*/NULL, + /*declarator_p=*/false, + /*optional_p=*/false); + decl = cp_parser_lookup_name_simple (parser, id, token->location); + if (id != error_mark_node && decl == error_mark_node) + cp_parser_name_lookup_error (parser, id, decl, NLE_NULL, + token->location); + + if (decl == error_mark_node + || !cp_parser_require (parser, CPP_CLOSE_PAREN, RT_CLOSE_PAREN)) + { + cp_parser_skip_to_pragma_eol (parser, pragma_tok); + return; + } + } + + /* Build a chain of clauses. */ + parser->lexer->in_pragma = true; + tree clauses = NULL_TREE; + clauses = cp_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK, + "#pragma acc routine", + cp_lexer_peek_token (parser->lexer)); + + /* Force clauses to be non-null, by attaching context to it. */ + clauses = tree_cons (c_head, clauses, NULL_TREE); + + if (decl) + cp_parser_finish_oacc_routine (parser, decl, clauses, true, false); + else + parser->oacc_routine = clauses; +} + +/* Apply any saved OpenACC routine clauses to a just-parsed + declaration. */ + +static void +cp_finalize_oacc_routine (cp_parser *parser, tree fndecl, bool is_defn) +{ + if (parser->oacc_routine) + { + cp_parser_finish_oacc_routine (parser, fndecl, parser->oacc_routine, + false, is_defn); + parser->oacc_routine = NULL_TREE; + } +} + /* Main entry point to OpenMP statement pragmas. */ static void @@ -36063,8 +36226,9 @@ cp_parser_pragma (cp_parser *parser, enu parser->lexer->in_pragma = true; id = pragma_tok->pragma_kind; - if (id != PRAGMA_OMP_DECLARE_REDUCTION) + if (id != PRAGMA_OMP_DECLARE_REDUCTION && id != PRAGMA_OACC_ROUTINE) cp_ensure_no_omp_declare_simd (parser); + cp_ensure_no_oacc_routine (parser); switch (id) { case PRAGMA_GCC_PCH_PREPROCESS: @@ -36174,6 +36338,10 @@ cp_parser_pragma (cp_parser *parser, enu cp_parser_omp_declare (parser, pragma_tok, context); return false; + case PRAGMA_OACC_ROUTINE: + cp_parser_oacc_routine (parser, pragma_tok, context); + return false; + case PRAGMA_OACC_ATOMIC: case PRAGMA_OACC_CACHE: case PRAGMA_OACC_DATA: Index: gcc/cp/parser.h =================================================================== --- gcc/cp/parser.h (revision 230040) +++ gcc/cp/parser.h (working copy) @@ -371,6 +371,9 @@ struct GTY(()) cp_parser { necessary. */ cp_omp_declare_simd_data * GTY((skip)) cilk_simd_fn_info; + /* OpenACC routine clauses for subsequent decl/defn. */ + tree oacc_routine; + /* Nonzero if parsing a parameter list where 'auto' should trigger an implicit template parameter. */ bool auto_is_implicit_function_template_parm_p; ^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [1/2] OpenACC routine support 2015-11-10 0:32 ` Nathan Sidwell @ 2015-11-10 0:48 ` Nathan Sidwell 2015-11-10 5:31 ` Cesar Philippidis 2015-11-10 5:28 ` Cesar Philippidis 1 sibling, 1 reply; 23+ messages in thread From: Nathan Sidwell @ 2015-11-10 0:48 UTC (permalink / raw) To: Jakub Jelinek; +Cc: GCC Patches, Cesar Philippidis [-- Attachment #1: Type: text/plain, Size: 167 bytes --] And these are the new tests. Cesar, c-c++-common/goacc/routine-5.c will need adjusting with your C++ parser patch. You'll see the two cases I've #if'd out. nathan [-- Attachment #2: trunk-routine-tests-1109.patch --] [-- Type: text/x-patch, Size: 14874 bytes --] 2015-11-09 Nathan Sidwell <nathan@codesourcery.com> gcc/testsuite/ * c-c++-common/goacc/routine-1.c: New. * c-c++-common/goacc/routine-2.c: New. * c-c++-common/goacc/routine-3.c: New. * c-c++-common/goacc/routine-4.c: New. * c-c++-common/goacc/routine-5.c: New. libgomp/ * testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c: New. * testsuite/libgomp.oacc-c-c++-common/routine-g-1.c: New. * testsuite/libgomp.oacc-c-c++-common/routine-gwv-1.c: New. * testsuite/libgomp.oacc-c-c++-common/routine-v-1.c: New. * testsuite/libgomp.oacc-c-c++-common/routine-w-1.c: New. * testsuite/libgomp.oacc-c-c++-common/routine-wv-1.c: New. Index: gcc/testsuite/c-c++-common/goacc/routine-1.c =================================================================== --- gcc/testsuite/c-c++-common/goacc/routine-1.c (revision 0) +++ gcc/testsuite/c-c++-common/goacc/routine-1.c (working copy) @@ -0,0 +1,34 @@ + +#pragma acc routine gang +void gang (void) +{ +} + +#pragma acc routine worker +void worker (void) +{ +} + +#pragma acc routine vector +void vector (void) +{ +} + +#pragma acc routine seq +void seq (void) +{ +} + +int main () +{ + +#pragma acc parallel num_gangs (32) num_workers (32) vector_length (32) + { + gang (); + worker (); + vector (); + seq (); + } + + return 0; +} Index: gcc/testsuite/c-c++-common/goacc/routine-2.c =================================================================== --- gcc/testsuite/c-c++-common/goacc/routine-2.c (revision 0) +++ gcc/testsuite/c-c++-common/goacc/routine-2.c (working copy) @@ -0,0 +1,21 @@ +#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) +{ +} + +#pragma acc routine (nothing) gang /* { dg-error "not been declared" } */ Index: gcc/testsuite/c-c++-common/goacc/routine-3.c =================================================================== --- gcc/testsuite/c-c++-common/goacc/routine-3.c (revision 0) +++ gcc/testsuite/c-c++-common/goacc/routine-3.c (working copy) @@ -0,0 +1,53 @@ +#pragma acc routine gang +void gang (void) /* { dg-message "declared here" 3 } */ +{ +} + +#pragma acc routine worker +void worker (void) /* { dg-message "declared here" 2 } */ +{ +} + +#pragma acc routine vector +void vector (void) /* { dg-message "declared here" 1 } */ +{ +} + +#pragma acc routine seq +void seq (void) +{ +} + +int main () +{ + +#pragma acc parallel num_gangs (32) num_workers (32) vector_length (32) + { + #pragma acc loop gang /* { dg-message "loop here" 1 } */ + for (int i = 0; i < 10; i++) + { + gang (); /* { dg-error "routine call uses same" } */ + worker (); + vector (); + seq (); + } + #pragma acc loop worker /* { dg-message "loop here" 2 } */ + for (int i = 0; i < 10; i++) + { + gang (); /* { dg-error "routine call uses same" } */ + worker (); /* { dg-error "routine call uses same" } */ + vector (); + seq (); + } + #pragma acc loop vector /* { dg-message "loop here" 3 } */ + for (int i = 0; i < 10; i++) + { + gang (); /* { dg-error "routine call uses same" } */ + worker (); /* { dg-error "routine call uses same" } */ + vector (); /* { dg-error "routine call uses same" } */ + seq (); + } + } + + return 0; +} Index: gcc/testsuite/c-c++-common/goacc/routine-4.c =================================================================== --- gcc/testsuite/c-c++-common/goacc/routine-4.c (revision 0) +++ gcc/testsuite/c-c++-common/goacc/routine-4.c (working copy) @@ -0,0 +1,41 @@ + +void gang (void); +void worker (void); +void vector (void); + +#pragma acc routine (gang) gang +#pragma acc routine (worker) worker +#pragma acc routine (vector) vector + +#pragma acc routine seq +void seq (void) +{ + gang (); /* { dg-error "routine call uses" } */ + worker (); /* { dg-error "routine call uses" } */ + vector (); /* { dg-error "routine call uses" } */ + seq (); +} + +void vector (void) /* { dg-message "declared here" 1 } */ +{ + gang (); /* { dg-error "routine call uses" } */ + worker (); /* { dg-error "routine call uses" } */ + vector (); + seq (); +} + +void worker (void) /* { dg-message "declared here" 2 } */ +{ + gang (); /* { dg-error "routine call uses" } */ + worker (); + vector (); + seq (); +} + +void gang (void) /* { dg-message "declared here" 3 } */ +{ + gang (); + worker (); + vector (); + seq (); +} Index: gcc/testsuite/c-c++-common/goacc/routine-5.c =================================================================== --- gcc/testsuite/c-c++-common/goacc/routine-5.c (revision 0) +++ gcc/testsuite/c-c++-common/goacc/routine-5.c (working copy) @@ -0,0 +1,51 @@ +/* { dg-do compile } */ + +#pragma acc routine /* { dg-error "not followed by" } */ +int a; + +#if 0 /* Disable for the moment. */ +#pragma acc routine /* dg-error "not followed by" */ +void fn1 (void), fn1b (void); +#endif + +#pragma acc routine /* { dg-error "not followed by" } */ +int b, fn2 (void); + +#if 0 /* Disable for the moment. */ +#pragma acc routine /* dg-error "not followed by" */ +int fn3 (void), b2; +#endif + +#pragma acc routine /* { dg-error "not followed by" } */ +typedef struct c c; + +#pragma acc routine /* { dg-error "not followed by" } */ +struct d {} d; + +#pragma acc routine /* { dg-error "not followed by" } */ +#pragma acc routine +int fn4 (void); + +int fn5a (void); + +#pragma acc routine /* { dg-error "not followed by" } */ +#pragma acc routine (fn5a) +int fn5 (void); + +#ifdef __cplusplus + +#pragma acc routine /* { dg-error "not followed by" "" { target c++ } } */ +namespace f {} + +namespace g {} + +#pragma acc routine /* { dg-error "not followed by" "" { target c++ } } */ +using namespace g; + +#pragma acc routine (g) /* { dg-error "does not refer to" "" { target c++ } } */ + +#endif + +#pragma acc routine (a) /* { dg-error "does not refer to" } */ + +#pragma acc routine (c) /* { dg-error "does not refer to" } */ Index: libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c (revision 0) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c (working copy) @@ -0,0 +1,41 @@ +/* { dg-do run } */ + +#include <openacc.h> + +int main () +{ + int ok = 1; + int val = 2; + int ary[32]; + int ondev = 0; + + for (int i = 0; i < 32; i++) + ary[i] = ~0; + +#pragma acc parallel num_gangs (32) copy (ok) firstprivate (val) copy(ary, ondev) + { + ondev = acc_on_device (acc_device_not_host); +#pragma acc loop gang(static:1) + for (unsigned i = 0; i < 32; i++) + { + if (val != 2) + ok = 0; + val += i; + ary[i] = val; + } + } + + if (ondev) + { + if (!ok) + return 1; + if (val != 2) + return 1; + + for (int i = 0; i < 32; i++) + if (ary[i] != 2 + i) + return 1; + } + + return 0; +} Index: libgomp/testsuite/libgomp.oacc-c-c++-common/routine-g-1.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/routine-g-1.c (revision 0) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/routine-g-1.c (working copy) @@ -0,0 +1,64 @@ +/* { dg-do run } */ +/* { dg-additional-options "-O2" */ + +#include <stdio.h> + +#define N (32*32*32+17) + +#pragma acc routine gang +void __attribute__ ((noinline)) gang (int ary[N]) +{ +#pragma acc loop gang + for (unsigned ix = 0; ix < N; ix++) + { + if (__builtin_acc_on_device (5)) + { + int g = 0, w = 0, v = 0; + + __asm__ volatile ("mov.u32 %0,%%ctaid.x;" : "=r" (g)); + __asm__ volatile ("mov.u32 %0,%%tid.y;" : "=r" (w)); + __asm__ volatile ("mov.u32 %0,%%tid.x;" : "=r" (v)); + ary[ix] = (g << 16) | (w << 8) | v; + } + else + ary[ix] = ix; + } +} + +int main () +{ + int ary[N]; + int ix; + int exit = 0; + int ondev = 0; + + for (ix = 0; ix < N;ix++) + ary[ix] = -1; + +#pragma acc parallel num_gangs(32) vector_length(32) copy(ary) copy(ondev) + { + ondev = __builtin_acc_on_device (5); + gang (ary); + } + + for (ix = 0; ix < N; ix++) + { + int expected = ix; + if(ondev) + { + int g = ix / ((N + 31) / 32); + int w = 0; + int v = 0; + + expected = (g << 16) | (w << 8) | v; + } + + if (ary[ix] != expected) + { + exit = 1; + printf ("ary[%d]=%x expected %x\n", ix, ary[ix], expected); + } + } + + return exit; +} Index: libgomp/testsuite/libgomp.oacc-c-c++-common/routine-gwv-1.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/routine-gwv-1.c (revision 0) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/routine-gwv-1.c (working copy) @@ -0,0 +1,66 @@ +/* { dg-do run } */ +/* { dg-additional-options "-O2" */ + +#include <stdio.h> + +#define N (32*32*32+17) + +#pragma acc routine gang +void __attribute__ ((noinline)) gang (int ary[N]) +{ +#pragma acc loop gang worker vector + for (unsigned ix = 0; ix < N; ix++) + { + if (__builtin_acc_on_device (5)) + { + int g = 0, w = 0, v = 0; + + __asm__ volatile ("mov.u32 %0,%%ctaid.x;" : "=r" (g)); + __asm__ volatile ("mov.u32 %0,%%tid.y;" : "=r" (w)); + __asm__ volatile ("mov.u32 %0,%%tid.x;" : "=r" (v)); + ary[ix] = (g << 16) | (w << 8) | v; + } + else + ary[ix] = ix; + } +} + +int main () +{ + int ary[N]; + int ix; + int exit = 0; + int ondev = 0; + + for (ix = 0; ix < N;ix++) + ary[ix] = -1; + +#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) copy(ary) copy(ondev) + { + ondev = __builtin_acc_on_device (5); + gang (ary); + } + + for (ix = 0; ix < N; ix++) + { + int expected = ix; + if(ondev) + { + int chunk_size = (N + 32*32*32 - 1) / (32*32*32); + + int g = ix / (chunk_size * 32 * 32); + int w = ix / 32 % 32; + int v = ix % 32; + + expected = (g << 16) | (w << 8) | v; + } + + if (ary[ix] != expected) + { + exit = 1; + printf ("ary[%d]=%x expected %x\n", ix, ary[ix], expected); + } + } + + return exit; +} Index: libgomp/testsuite/libgomp.oacc-c-c++-common/routine-v-1.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/routine-v-1.c (revision 0) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/routine-v-1.c (working copy) @@ -0,0 +1,64 @@ +/* { dg-do run } */ +/* { dg-additional-options "-O2" */ + +#include <stdio.h> + +#define N (32*32*32+17) + +#pragma acc routine vector +void __attribute__ ((noinline)) vector (int ary[N]) +{ +#pragma acc loop vector + for (unsigned ix = 0; ix < N; ix++) + { + if (__builtin_acc_on_device (5)) + { + int g = 0, w = 0, v = 0; + + __asm__ volatile ("mov.u32 %0,%%ctaid.x;" : "=r" (g)); + __asm__ volatile ("mov.u32 %0,%%tid.y;" : "=r" (w)); + __asm__ volatile ("mov.u32 %0,%%tid.x;" : "=r" (v)); + ary[ix] = (g << 16) | (w << 8) | v; + } + else + ary[ix] = ix; + } +} + +int main () +{ + int ary[N]; + int ix; + int exit = 0; + int ondev = 0; + + for (ix = 0; ix < N;ix++) + ary[ix] = -1; + +#pragma acc parallel vector_length(32) copy(ary) copy(ondev) + { + ondev = __builtin_acc_on_device (5); + vector (ary); + } + + for (ix = 0; ix < N; ix++) + { + int expected = ix; + if(ondev) + { + int g = 0; + int w = 0; + int v = ix % 32; + + expected = (g << 16) | (w << 8) | v; + } + + if (ary[ix] != expected) + { + exit = 1; + printf ("ary[%d]=%x expected %x\n", ix, ary[ix], expected); + } + } + + return exit; +} Index: libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c (revision 0) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c (working copy) @@ -0,0 +1,64 @@ +/* { dg-do run } */ +/* { dg-additional-options "-O2" */ + +#include <stdio.h> + +#define N (32*32*32+17) + +#pragma acc routine worker +void __attribute__ ((noinline)) worker (int ary[N]) +{ +#pragma acc loop worker + for (unsigned ix = 0; ix < N; ix++) + { + if (__builtin_acc_on_device (5)) + { + int g = 0, w = 0, v = 0; + + __asm__ volatile ("mov.u32 %0,%%ctaid.x;" : "=r" (g)); + __asm__ volatile ("mov.u32 %0,%%tid.y;" : "=r" (w)); + __asm__ volatile ("mov.u32 %0,%%tid.x;" : "=r" (v)); + ary[ix] = (g << 16) | (w << 8) | v; + } + else + ary[ix] = ix; + } +} + +int main () +{ + int ary[N]; + int ix; + int exit = 0; + int ondev = 0; + + for (ix = 0; ix < N;ix++) + ary[ix] = -1; + +#pragma acc parallel num_workers(32) vector_length(32) copy(ary) copy(ondev) + { + ondev = __builtin_acc_on_device (5); + worker (ary); + } + + for (ix = 0; ix < N; ix++) + { + int expected = ix; + if(ondev) + { + int g = 0; + int w = ix % 32; + int v = 0; + + expected = (g << 16) | (w << 8) | v; + } + + if (ary[ix] != expected) + { + exit = 1; + printf ("ary[%d]=%x expected %x\n", ix, ary[ix], expected); + } + } + + return exit; +} Index: libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-1.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-1.c (revision 0) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-1.c (working copy) @@ -0,0 +1,64 @@ +/* { dg-do run } */ +/* { dg-additional-options "-O2" */ + +#include <stdio.h> + +#define N (32*32*32+17) + +#pragma acc routine worker +void __attribute__ ((noinline)) worker (int ary[N]) +{ +#pragma acc loop worker vector + for (unsigned ix = 0; ix < N; ix++) + { + if (__builtin_acc_on_device (5)) + { + int g = 0, w = 0, v = 0; + + __asm__ volatile ("mov.u32 %0,%%ctaid.x;" : "=r" (g)); + __asm__ volatile ("mov.u32 %0,%%tid.y;" : "=r" (w)); + __asm__ volatile ("mov.u32 %0,%%tid.x;" : "=r" (v)); + ary[ix] = (g << 16) | (w << 8) | v; + } + else + ary[ix] = ix; + } +} + +int main () +{ + int ary[N]; + int ix; + int exit = 0; + int ondev = 0; + + for (ix = 0; ix < N;ix++) + ary[ix] = -1; + +#pragma acc parallel num_workers(32) vector_length(32) copy(ary) copy(ondev) + { + ondev = __builtin_acc_on_device (5); + worker (ary); + } + + for (ix = 0; ix < N; ix++) + { + int expected = ix; + if(ondev) + { + int g = 0; + int w = (ix / 32) % 32; + int v = ix % 32; + + expected = (g << 16) | (w << 8) | v; + } + + if (ary[ix] != expected) + { + exit = 1; + printf ("ary[%d]=%x expected %x\n", ix, ary[ix], expected); + } + } + + return exit; +} ^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [1/2] OpenACC routine support 2015-11-10 0:48 ` Nathan Sidwell @ 2015-11-10 5:31 ` Cesar Philippidis 0 siblings, 0 replies; 23+ messages in thread From: Cesar Philippidis @ 2015-11-10 5:31 UTC (permalink / raw) To: Nathan Sidwell, Jakub Jelinek; +Cc: GCC Patches On 11/09/2015 04:48 PM, Nathan Sidwell wrote: > And these are the new tests. Cesar, c-c++-common/goacc/routine-5.c will > need adjusting with your C++ parser patch. You'll see the two cases > I've #if'd out. I enabled those tests in trunk with the patch I posted here <https://gcc.gnu.org/ml/gcc-patches/2015-11/msg01150.html>. Cesar ^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [1/2] OpenACC routine support 2015-11-10 0:32 ` Nathan Sidwell 2015-11-10 0:48 ` Nathan Sidwell @ 2015-11-10 5:28 ` Cesar Philippidis 2015-11-10 8:16 ` Jakub Jelinek 1 sibling, 1 reply; 23+ messages in thread From: Cesar Philippidis @ 2015-11-10 5:28 UTC (permalink / raw) To: Nathan Sidwell, Jakub Jelinek; +Cc: GCC Patches [-- Attachment #1: Type: text/plain, Size: 1579 bytes --] On 11/09/2015 04:31 PM, Nathan Sidwell wrote: > On 11/03/15 10:35, Jakub Jelinek wrote: >> On Mon, Nov 02, 2015 at 02:21:43PM -0500, Nathan Sidwell wrote: >>> --- gcc/c/c-parser.c (revision 229667) >>> +++ gcc/c/c-parser.c (working copy) >>> @@ -1160,7 +1160,8 @@ enum c_parser_prec { >>> static void c_parser_external_declaration (c_parser *); >>> static void c_parser_asm_definition (c_parser *); >>> static void c_parser_declaration_or_fndef (c_parser *, bool, bool, >>> bool, >>> - bool, bool, tree *, vec<c_token>); >>> + bool, bool, tree *, vec<c_token>, >>> + tree); >> >> Wonder if this shouldn't be tree = NULL_TREE, then you'd avoid most of >> the >> c_parser_declaration_or_fndef caller changes. >> >> Otherwise, LGTM. > > This is the patch I've just committed. It includes c parser adjustments > to detect the case of two function decls with a single type specifier. > Cesar will be applying a patch for the C++ parser for the same case. Here's the patch that Nathan was referring to. I ended up introducing a boolean variable named first in the various functions which call finalize_oacc_routines. The problem the original approach was having was that the routine clauses is only applied to the first function declarator in a declaration list. By using 'first', which is set to true if the current declarator is the first in a sequence of declarators, I was able to defer setting parser->oacc_routine to NULL. Nathan already approved this patch, so I've applied it to trunk. Cesar [-- Attachment #2: oacc-routines-trunk.diff --] [-- Type: text/x-patch, Size: 11588 bytes --] 2015-11-09 Cesar Philippidis <cesar@codesourcery.com> gcc/cp/ * parser.c (cp_finalize_oacc_routine): New boolean first argument. (cp_ensure_no_oacc_routine): Update call to cp_finalize_oacc_routine. (cp_parser_simple_declaration): Maintain a boolean first to keep track of each new declarator. Propagate it to cp_parser_init_declarator. (cp_parser_init_declarator): New boolean first argument. Propagate it to cp_parser_save_member_function_body and cp_finalize_oacc_routine. (cp_parser_member_declaration): Likewise. (cp_parser_single_declaration): Update call to cp_parser_init_declarator. (cp_parser_save_member_function_body): New boolean first_decl argument. Propagate it to cp_finalize_oacc_routine. (cp_parser_finish_oacc_routine): New boolean first argument. Use it to determine if multiple declarators follow a routine construct. (cp_parser_oacc_routine): Update call to cp_parser_finish_oacc_routine. gcc/testsuite/ * c-c++-common/goacc/routine-5.c: Enable c++ tests. diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index 6fc2c6a..f3b4b46 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -246,7 +246,7 @@ static bool cp_parser_omp_declare_reduction_exprs static tree cp_parser_cilk_simd_vectorlength (cp_parser *, tree, bool); static void cp_finalize_oacc_routine - (cp_parser *, tree, bool); + (cp_parser *, tree, bool, bool); /* Manifest constants. */ #define CP_LEXER_BUFFER_SIZE ((256 * 1024) / sizeof (cp_token)) @@ -1329,7 +1329,7 @@ cp_finalize_omp_declare_simd (cp_parser *parser, tree fndecl) static inline void cp_ensure_no_oacc_routine (cp_parser *parser) { - cp_finalize_oacc_routine (parser, NULL_TREE, false); + cp_finalize_oacc_routine (parser, NULL_TREE, false, true); } \f /* Decl-specifiers. */ @@ -2135,7 +2135,7 @@ static tree cp_parser_decltype static tree cp_parser_init_declarator (cp_parser *, cp_decl_specifier_seq *, vec<deferred_access_check, va_gc> *, - bool, bool, int, bool *, tree *, location_t *); + bool, bool, int, bool *, tree *, bool, location_t *); static cp_declarator *cp_parser_declarator (cp_parser *, cp_parser_declarator_kind, int *, bool *, bool, bool); static cp_declarator *cp_parser_direct_declarator @@ -2445,7 +2445,7 @@ static tree cp_parser_single_declaration static tree cp_parser_functional_cast (cp_parser *, tree); static tree cp_parser_save_member_function_body - (cp_parser *, cp_decl_specifier_seq *, cp_declarator *, tree); + (cp_parser *, cp_decl_specifier_seq *, cp_declarator *, tree, bool); static tree cp_parser_save_nsdmi (cp_parser *); static tree cp_parser_enclosed_template_argument_list @@ -11909,6 +11909,7 @@ cp_parser_simple_declaration (cp_parser* parser, bool saw_declarator; location_t comma_loc = UNKNOWN_LOCATION; location_t init_loc = UNKNOWN_LOCATION; + bool first = true; if (maybe_range_for_decl) *maybe_range_for_decl = NULL_TREE; @@ -12005,7 +12006,10 @@ cp_parser_simple_declaration (cp_parser* parser, declares_class_or_enum, &function_definition_p, maybe_range_for_decl, + first, &init_loc); + first = false; + /* If an error occurred while parsing tentatively, exit quickly. (That usually happens when in the body of a function; each statement is treated as a declaration-statement until proven @@ -12104,6 +12108,9 @@ cp_parser_simple_declaration (cp_parser* parser, done: pop_deferring_access_checks (); + + /* Reset any acc routine clauses. */ + parser->oacc_routine = NULL; } /* Parse a decl-specifier-seq. @@ -17843,6 +17850,8 @@ cp_parser_asm_definition (cp_parser* parser) if present, will not be consumed. If returned, this declarator will be created with SD_INITIALIZED but will not call cp_finish_decl. + FIRST indicates if this is the first declarator in a declaration sequence. + If INIT_LOC is not NULL, and *INIT_LOC is equal to UNKNOWN_LOCATION, and there is an initializer, the pointed location_t is set to the location of the '=' or `(', or '{' in C++11 token introducing the @@ -17857,6 +17866,7 @@ cp_parser_init_declarator (cp_parser* parser, int declares_class_or_enum, bool* function_definition_p, tree* maybe_range_for_decl, + bool first, location_t* init_loc) { cp_token *token = NULL, *asm_spec_start_token = NULL, @@ -17993,7 +18003,8 @@ cp_parser_init_declarator (cp_parser* parser, decl = cp_parser_save_member_function_body (parser, decl_specifiers, declarator, - prefix_attributes); + prefix_attributes, + true); else decl = (cp_parser_function_definition_from_specifiers_and_declarator @@ -18097,7 +18108,7 @@ cp_parser_init_declarator (cp_parser* parser, range_for_decl_p? SD_INITIALIZED : is_initialized, attributes, prefix_attributes, &pushed_scope); cp_finalize_omp_declare_simd (parser, decl); - cp_finalize_oacc_routine (parser, decl, false); + cp_finalize_oacc_routine (parser, decl, false, first); /* Adjust location of decl if declarator->id_loc is more appropriate: set, and decl wasn't merged with another decl, in which case its location would be different from input_location, and more accurate. */ @@ -18211,7 +18222,7 @@ cp_parser_init_declarator (cp_parser* parser, if (decl && TREE_CODE (decl) == FUNCTION_DECL) cp_parser_save_default_args (parser, decl); cp_finalize_omp_declare_simd (parser, decl); - cp_finalize_oacc_routine (parser, decl, false); + cp_finalize_oacc_routine (parser, decl, false, first); } /* Finish processing the declaration. But, skip member @@ -21915,6 +21926,7 @@ cp_parser_member_declaration (cp_parser* parser) else { bool assume_semicolon = false; + bool first = true; /* Clear attributes from the decl_specifiers but keep them around as prefix attributes that apply them to the entity @@ -22102,7 +22114,10 @@ cp_parser_member_declaration (cp_parser* parser) decl = cp_parser_save_member_function_body (parser, &decl_specifiers, declarator, - attributes); + attributes, + first); + first = false; + if (parser->fully_implicit_function_template_p) decl = finish_fully_implicit_template (parser, decl); /* If the member was not a friend, declare it here. */ @@ -22132,7 +22147,8 @@ cp_parser_member_declaration (cp_parser* parser) } cp_finalize_omp_declare_simd (parser, decl); - cp_finalize_oacc_routine (parser, decl, false); + cp_finalize_oacc_routine (parser, decl, false, first); + first = false; /* Reset PREFIX_ATTRIBUTES. */ while (attributes && TREE_CHAIN (attributes) != first_attribute) @@ -22195,6 +22211,9 @@ cp_parser_member_declaration (cp_parser* parser) if (assume_semicolon) goto out; } + + /* Reset any OpenACC routine clauses. */ + parser->oacc_routine = NULL; } cp_parser_require (parser, CPP_SEMICOLON, RT_SEMICOLON); @@ -24736,7 +24755,8 @@ cp_parser_function_definition_from_specifiers_and_declarator { cp_finalize_omp_declare_simd (parser, current_function_decl); parser->omp_declare_simd = NULL; - cp_finalize_oacc_routine (parser, current_function_decl, true); + cp_finalize_oacc_routine (parser, current_function_decl, true, true); + parser->oacc_routine = NULL; } if (!success_p) @@ -25301,7 +25321,7 @@ cp_parser_single_declaration (cp_parser* parser, member_p, declares_class_or_enum, &function_definition_p, - NULL, NULL); + NULL, true, NULL); /* 7.1.1-1 [dcl.stc] @@ -25403,14 +25423,15 @@ cp_parser_functional_cast (cp_parser* parser, tree type) /* Save the tokens that make up the body of a member function defined in a class-specifier. The DECL_SPECIFIERS and DECLARATOR have already been parsed. The ATTRIBUTES are any GNU "__attribute__" - specifiers applied to the declaration. Returns the FUNCTION_DECL - for the member function. */ + specifiers applied to the declaration. FIRST_DECL indicates if + DECLARATOR is the first declarator in a declaration sequence. Returns + the FUNCTION_DECL for the member function. */ static tree cp_parser_save_member_function_body (cp_parser* parser, cp_decl_specifier_seq *decl_specifiers, cp_declarator *declarator, - tree attributes) + tree attributes, bool first_decl) { cp_token *first; cp_token *last; @@ -25419,7 +25440,7 @@ cp_parser_save_member_function_body (cp_parser* parser, /* Create the FUNCTION_DECL. */ fn = grokmethod (decl_specifiers, declarator, attributes); cp_finalize_omp_declare_simd (parser, fn); - cp_finalize_oacc_routine (parser, fn, true); + cp_finalize_oacc_routine (parser, fn, true, first_decl); /* If something went badly wrong, bail out now. */ if (fn == error_mark_node) { @@ -35625,7 +35646,8 @@ cp_parser_omp_taskloop (cp_parser *parser, cp_token *pragma_tok, static void cp_parser_finish_oacc_routine (cp_parser *ARG_UNUSED (parser), tree fndecl, - tree clauses, bool named, bool is_defn) + tree clauses, bool named, bool is_defn, + bool first) { location_t loc = OMP_CLAUSE_LOCATION (TREE_PURPOSE (clauses)); @@ -35637,7 +35659,8 @@ cp_parser_finish_oacc_routine (cp_parser *ARG_UNUSED (parser), tree fndecl, return; } - if (!fndecl || TREE_CODE (fndecl) != FUNCTION_DECL) + if (!fndecl || TREE_CODE (fndecl) != FUNCTION_DECL + || (!named && !first)) { error_at (loc, "%<#pragma acc routine%> %s", named ? "does not refer to a function" @@ -35728,7 +35751,7 @@ cp_parser_oacc_routine (cp_parser *parser, cp_token *pragma_tok, clauses = tree_cons (c_head, clauses, NULL_TREE); if (decl) - cp_parser_finish_oacc_routine (parser, decl, clauses, true, false); + cp_parser_finish_oacc_routine (parser, decl, clauses, true, false, 0); else parser->oacc_routine = clauses; } @@ -35737,14 +35760,12 @@ cp_parser_oacc_routine (cp_parser *parser, cp_token *pragma_tok, declaration. */ static void -cp_finalize_oacc_routine (cp_parser *parser, tree fndecl, bool is_defn) +cp_finalize_oacc_routine (cp_parser *parser, tree fndecl, bool is_defn, + bool first) { if (parser->oacc_routine) - { - cp_parser_finish_oacc_routine (parser, fndecl, parser->oacc_routine, - false, is_defn); - parser->oacc_routine = NULL_TREE; - } + cp_parser_finish_oacc_routine (parser, fndecl, parser->oacc_routine, + false, is_defn, first); } /* Main entry point to OpenMP statement pragmas. */ diff --git a/gcc/testsuite/c-c++-common/goacc/routine-5.c b/gcc/testsuite/c-c++-common/goacc/routine-5.c index c2a8fb2..ccda097 100644 --- a/gcc/testsuite/c-c++-common/goacc/routine-5.c +++ b/gcc/testsuite/c-c++-common/goacc/routine-5.c @@ -3,18 +3,14 @@ #pragma acc routine /* { dg-error "not followed by" } */ int a; -#if 0 /* Disable for the moment. */ -#pragma acc routine /* dg-error "not followed by" */ +#pragma acc routine /* { dg-error "not followed by" } */ void fn1 (void), fn1b (void); -#endif #pragma acc routine /* { dg-error "not followed by" } */ int b, fn2 (void); -#if 0 /* Disable for the moment. */ -#pragma acc routine /* dg-error "not followed by" */ +#pragma acc routine /* { dg-error "not followed by" } */ int fn3 (void), b2; -#endif #pragma acc routine /* { dg-error "not followed by" } */ typedef struct c c; ^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [1/2] OpenACC routine support 2015-11-10 5:28 ` Cesar Philippidis @ 2015-11-10 8:16 ` Jakub Jelinek 2015-11-10 14:37 ` Cesar Philippidis 2015-11-18 19:02 ` Cesar Philippidis 0 siblings, 2 replies; 23+ messages in thread From: Jakub Jelinek @ 2015-11-10 8:16 UTC (permalink / raw) To: Cesar Philippidis; +Cc: Nathan Sidwell, GCC Patches On Mon, Nov 09, 2015 at 09:28:47PM -0800, Cesar Philippidis wrote: > Here's the patch that Nathan was referring to. I ended up introducing a > boolean variable named first in the various functions which call > finalize_oacc_routines. The problem the original approach was having was > that the routine clauses is only applied to the first function > declarator in a declaration list. By using 'first', which is set to true > if the current declarator is the first in a sequence of declarators, I > was able to defer setting parser->oacc_routine to NULL. The #pragma omp declare simd has identical restrictions, but doesn't need to add any of the first parameters to the C++ parser. So, what are you doing differently that you need it? Handling both differently is a consistency issue, and unnecessary additional complexity to the parser. Jakub ^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [1/2] OpenACC routine support 2015-11-10 8:16 ` Jakub Jelinek @ 2015-11-10 14:37 ` Cesar Philippidis 2015-11-18 19:02 ` Cesar Philippidis 1 sibling, 0 replies; 23+ messages in thread From: Cesar Philippidis @ 2015-11-10 14:37 UTC (permalink / raw) To: Jakub Jelinek; +Cc: Nathan Sidwell, GCC Patches On 11/10/2015 12:16 AM, Jakub Jelinek wrote: > On Mon, Nov 09, 2015 at 09:28:47PM -0800, Cesar Philippidis wrote: >> Here's the patch that Nathan was referring to. I ended up introducing a >> boolean variable named first in the various functions which call >> finalize_oacc_routines. The problem the original approach was having was >> that the routine clauses is only applied to the first function >> declarator in a declaration list. By using 'first', which is set to true >> if the current declarator is the first in a sequence of declarators, I >> was able to defer setting parser->oacc_routine to NULL. > > The #pragma omp declare simd has identical restrictions, but doesn't need > to add any of the first parameters to the C++ parser. > So, what are you doing differently that you need it? Handling both > differently is a consistency issue, and unnecessary additional complexity to > the parser. I see that you added an omp_declare_simd->fndecl_seen field to cp_parser. My objective was to try and make the c++ routine parsing somewhat consistent with the c front end. I could probably add a similar oacc_routine field, but I wonder if it would be better to share omp_declare_simd. There was talk about the next version of openacc adding support for -fopenacc and -fopenmp together. So maybe there needs to be a separate oacc_routine field. Cesar ^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [1/2] OpenACC routine support 2015-11-10 8:16 ` Jakub Jelinek 2015-11-10 14:37 ` Cesar Philippidis @ 2015-11-18 19:02 ` Cesar Philippidis 2015-11-19 11:23 ` Jakub Jelinek 2015-12-01 14:40 ` Thomas Schwinge 1 sibling, 2 replies; 23+ messages in thread From: Cesar Philippidis @ 2015-11-18 19:02 UTC (permalink / raw) To: Jakub Jelinek; +Cc: Nathan Sidwell, GCC Patches [-- Attachment #1: Type: text/plain, Size: 2008 bytes --] On 11/10/2015 12:16 AM, Jakub Jelinek wrote: > On Mon, Nov 09, 2015 at 09:28:47PM -0800, Cesar Philippidis wrote: >> Here's the patch that Nathan was referring to. I ended up introducing a >> boolean variable named first in the various functions which call >> finalize_oacc_routines. The problem the original approach was having was >> that the routine clauses is only applied to the first function >> declarator in a declaration list. By using 'first', which is set to true >> if the current declarator is the first in a sequence of declarators, I >> was able to defer setting parser->oacc_routine to NULL. > > The #pragma omp declare simd has identical restrictions, but doesn't need > to add any of the first parameters to the C++ parser. > So, what are you doing differently that you need it? Handling both > differently is a consistency issue, and unnecessary additional complexity to > the parser. I reworked how acc routines are handed in this patch to be more similar to #pragma omp declare simd. Things get kind of messy though. For starters, I had to add a new tree clauses member to cp_omp_declare_simd_data. This serves two purposes: * It allows the c++ FE to record the location of the first #pragma acc routine, which is nice because it allows test cases to be shared with the c FE. * Unlike omp declare simd, only one acc routine may be associated with a function decl. This meant that I had to defer attaching the acc geometry and 'omp target' attributes to cp_finalize_oacc_routine instead of in cp_parser_late_parsing_oacc_routine like in omp. So what happens is, cp_parser_late_parsing_oacc_routine ends up creating a function geometry clause. I don't really like this approach. I did try to postpone parsing the clauses till cp_finalize_oacc_routine, but that got messy. Plus, while I'd be able to remove the clauses field from cp_omp_declare_simd_data, we'd still need a location_t field for cp_ensure_no_oacc_routine. Is this OK for trunk? Cesar [-- Attachment #2: cxx-routine-v2.diff --] [-- Type: text/x-patch, Size: 22989 bytes --] 2015-11-17 Cesar Philippidis <cesar@codesourcery.com> gcc/cp/ * parser.h (struct cp_omp_declare_simd_data): Add clauses member. (struct cp_parser): Change type the of oacc_routine to cp_omp_declare_simd_data. * parser.c (cp_ensure_no_oacc_routine): Rework to use cp_omp_declare_simd_data. (cp_parser_simple_declaration): Remove boolean first. Update call to cp_parser_init_declarator. Don't NULL out oacc_routine. (cp_parser_init_declarator): Remove boolean first parameter. Update calls to cp_finalize_oacc_routine. (cp_parser_late_return_type_opt): Handle acc routines. (cp_parser_member_declaration): Remove first variable. Handle acc routines like omp declare simd. (cp_parser_function_definition_from_specifiers_and_declarator): Update call to cp_finalize_oacc_routine. (cp_parser_single_declaration): Update call to cp_parser_init_declarator. (cp_parser_save_member_function_body): Remove first_decl parameter. Update call to cp_finalize_oacc_routine. (cp_parser_finish_oacc_routine): Delete. (cp_parser_oacc_routine): Rework to use cp_omp_declare_simd_data. (cp_parser_late_parsing_oacc_routine): New function. (cp_finalize_oacc_routine): Remove first argument. Add more error handling and set the acc routine and 'omp declare target' attributes. (cp_parser_pragma): Remove unnecessary call to cp_ensure_no_oacc_routine. diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index 0e1116b..8de3bce 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -241,7 +241,7 @@ static bool cp_parser_omp_declare_reduction_exprs static tree cp_parser_cilk_simd_vectorlength (cp_parser *, tree, bool); static void cp_finalize_oacc_routine - (cp_parser *, tree, bool, bool); + (cp_parser *, tree, bool); /* Manifest constants. */ #define CP_LEXER_BUFFER_SIZE ((256 * 1024) / sizeof (cp_token)) @@ -1318,13 +1318,21 @@ cp_finalize_omp_declare_simd (cp_parser *parser, tree fndecl) } } -/* Diagnose if #pragma omp routine isn't followed immediately - by function declaration or definition. */ +/* Diagnose if #pragma acc routine isn't followed immediately by function + declaration or definition. */ static inline void cp_ensure_no_oacc_routine (cp_parser *parser) { - cp_finalize_oacc_routine (parser, NULL_TREE, false, true); + if (parser->oacc_routine && !parser->oacc_routine->error_seen) + { + tree clauses = parser->oacc_routine->clauses; + location_t loc = OMP_CLAUSE_LOCATION (TREE_PURPOSE(clauses)); + + error_at (loc, "%<#pragma oacc routine%> not followed by function " + "declaration or definition"); + parser->oacc_routine = NULL; + } } \f /* Decl-specifiers. */ @@ -2130,7 +2138,7 @@ static tree cp_parser_decltype static tree cp_parser_init_declarator (cp_parser *, cp_decl_specifier_seq *, vec<deferred_access_check, va_gc> *, - bool, bool, int, bool *, tree *, bool, location_t *); + bool, bool, int, bool *, tree *, location_t *); static cp_declarator *cp_parser_declarator (cp_parser *, cp_parser_declarator_kind, int *, bool *, bool, bool); static cp_declarator *cp_parser_direct_declarator @@ -2186,6 +2194,9 @@ static tree cp_parser_late_parsing_omp_declare_simd static tree cp_parser_late_parsing_cilk_simd_fn_info (cp_parser *, tree); +static tree cp_parser_late_parsing_oacc_routine + (cp_parser *, tree); + static tree synthesize_implicit_template_parm (cp_parser *, tree); static tree finish_fully_implicit_template @@ -2440,7 +2451,7 @@ static tree cp_parser_single_declaration static tree cp_parser_functional_cast (cp_parser *, tree); static tree cp_parser_save_member_function_body - (cp_parser *, cp_decl_specifier_seq *, cp_declarator *, tree, bool); + (cp_parser *, cp_decl_specifier_seq *, cp_declarator *, tree); static tree cp_parser_save_nsdmi (cp_parser *); static tree cp_parser_enclosed_template_argument_list @@ -11870,7 +11881,6 @@ cp_parser_simple_declaration (cp_parser* parser, bool saw_declarator; location_t comma_loc = UNKNOWN_LOCATION; location_t init_loc = UNKNOWN_LOCATION; - bool first = true; if (maybe_range_for_decl) *maybe_range_for_decl = NULL_TREE; @@ -11967,10 +11977,7 @@ cp_parser_simple_declaration (cp_parser* parser, declares_class_or_enum, &function_definition_p, maybe_range_for_decl, - first, &init_loc); - first = false; - /* If an error occurred while parsing tentatively, exit quickly. (That usually happens when in the body of a function; each statement is treated as a declaration-statement until proven @@ -12069,9 +12076,6 @@ cp_parser_simple_declaration (cp_parser* parser, done: pop_deferring_access_checks (); - - /* Reset any acc routine clauses. */ - parser->oacc_routine = NULL; } /* Parse a decl-specifier-seq. @@ -17811,8 +17815,6 @@ cp_parser_asm_definition (cp_parser* parser) if present, will not be consumed. If returned, this declarator will be created with SD_INITIALIZED but will not call cp_finish_decl. - FIRST indicates if this is the first declarator in a declaration sequence. - If INIT_LOC is not NULL, and *INIT_LOC is equal to UNKNOWN_LOCATION, and there is an initializer, the pointed location_t is set to the location of the '=' or `(', or '{' in C++11 token introducing the @@ -17827,7 +17829,6 @@ cp_parser_init_declarator (cp_parser* parser, int declares_class_or_enum, bool* function_definition_p, tree* maybe_range_for_decl, - bool first, location_t* init_loc) { cp_token *token = NULL, *asm_spec_start_token = NULL, @@ -17964,8 +17965,7 @@ cp_parser_init_declarator (cp_parser* parser, decl = cp_parser_save_member_function_body (parser, decl_specifiers, declarator, - prefix_attributes, - true); + prefix_attributes); else decl = (cp_parser_function_definition_from_specifiers_and_declarator @@ -18069,7 +18069,7 @@ cp_parser_init_declarator (cp_parser* parser, range_for_decl_p? SD_INITIALIZED : is_initialized, attributes, prefix_attributes, &pushed_scope); cp_finalize_omp_declare_simd (parser, decl); - cp_finalize_oacc_routine (parser, decl, false, first); + cp_finalize_oacc_routine (parser, decl, false); /* Adjust location of decl if declarator->id_loc is more appropriate: set, and decl wasn't merged with another decl, in which case its location would be different from input_location, and more accurate. */ @@ -18183,7 +18183,7 @@ cp_parser_init_declarator (cp_parser* parser, if (decl && TREE_CODE (decl) == FUNCTION_DECL) cp_parser_save_default_args (parser, decl); cp_finalize_omp_declare_simd (parser, decl); - cp_finalize_oacc_routine (parser, decl, false, first); + cp_finalize_oacc_routine (parser, decl, false); } /* Finish processing the declaration. But, skip member @@ -19289,13 +19289,17 @@ cp_parser_late_return_type_opt (cp_parser* parser, cp_declarator *declarator, bool cilk_simd_fn_vector_p = (parser->cilk_simd_fn_info && declarator && declarator->kind == cdk_id); - + + bool oacc_routine_p = (parser->oacc_routine + && declarator + && declarator->kind == cdk_id); + /* Peek at the next token. */ token = cp_lexer_peek_token (parser->lexer); /* A late-specified return type is indicated by an initial '->'. */ if (token->type != CPP_DEREF && token->keyword != RID_REQUIRES - && !(declare_simd_p || cilk_simd_fn_vector_p)) + && !(declare_simd_p || cilk_simd_fn_vector_p || oacc_routine_p)) return NULL_TREE; tree save_ccp = current_class_ptr; @@ -19326,7 +19330,11 @@ cp_parser_late_return_type_opt (cp_parser* parser, cp_declarator *declarator, declarator->std_attributes = cp_parser_late_parsing_omp_declare_simd (parser, declarator->std_attributes); - + if (oacc_routine_p) + declarator->std_attributes + = cp_parser_late_parsing_oacc_routine (parser, + declarator->std_attributes); + if (quals >= 0) { current_class_ptr = save_ccp; @@ -21887,7 +21895,6 @@ cp_parser_member_declaration (cp_parser* parser) else { bool assume_semicolon = false; - bool first = true; /* Clear attributes from the decl_specifiers but keep them around as prefix attributes that apply them to the entity @@ -22075,10 +22082,7 @@ cp_parser_member_declaration (cp_parser* parser) decl = cp_parser_save_member_function_body (parser, &decl_specifiers, declarator, - attributes, - first); - first = false; - + attributes); if (parser->fully_implicit_function_template_p) decl = finish_fully_implicit_template (parser, decl); /* If the member was not a friend, declare it here. */ @@ -22108,8 +22112,7 @@ cp_parser_member_declaration (cp_parser* parser) } cp_finalize_omp_declare_simd (parser, decl); - cp_finalize_oacc_routine (parser, decl, false, first); - first = false; + cp_finalize_oacc_routine (parser, decl, false); /* Reset PREFIX_ATTRIBUTES. */ while (attributes && TREE_CHAIN (attributes) != first_attribute) @@ -22172,9 +22175,6 @@ cp_parser_member_declaration (cp_parser* parser) if (assume_semicolon) goto out; } - - /* Reset any OpenACC routine clauses. */ - parser->oacc_routine = NULL; } cp_parser_require (parser, CPP_SEMICOLON, RT_SEMICOLON); @@ -24716,7 +24716,7 @@ cp_parser_function_definition_from_specifiers_and_declarator { cp_finalize_omp_declare_simd (parser, current_function_decl); parser->omp_declare_simd = NULL; - cp_finalize_oacc_routine (parser, current_function_decl, true, true); + cp_finalize_oacc_routine (parser, current_function_decl, true); parser->oacc_routine = NULL; } @@ -25282,7 +25282,7 @@ cp_parser_single_declaration (cp_parser* parser, member_p, declares_class_or_enum, &function_definition_p, - NULL, true, NULL); + NULL, NULL); /* 7.1.1-1 [dcl.stc] @@ -25384,15 +25384,14 @@ cp_parser_functional_cast (cp_parser* parser, tree type) /* Save the tokens that make up the body of a member function defined in a class-specifier. The DECL_SPECIFIERS and DECLARATOR have already been parsed. The ATTRIBUTES are any GNU "__attribute__" - specifiers applied to the declaration. FIRST_DECL indicates if - DECLARATOR is the first declarator in a declaration sequence. Returns - the FUNCTION_DECL for the member function. */ + specifiers applied to the declaration. Returns the FUNCTION_DECL + for the member function. */ static tree cp_parser_save_member_function_body (cp_parser* parser, cp_decl_specifier_seq *decl_specifiers, cp_declarator *declarator, - tree attributes, bool first_decl) + tree attributes) { cp_token *first; cp_token *last; @@ -25401,7 +25400,7 @@ cp_parser_save_member_function_body (cp_parser* parser, /* Create the FUNCTION_DECL. */ fn = grokmethod (decl_specifiers, declarator, attributes); cp_finalize_omp_declare_simd (parser, fn); - cp_finalize_oacc_routine (parser, fn, true, first_decl); + cp_finalize_oacc_routine (parser, fn, true); /* If something went badly wrong, bail out now. */ if (fn == error_mark_node) { @@ -35773,59 +35772,6 @@ cp_parser_omp_taskloop (cp_parser *parser, cp_token *pragma_tok, | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ)) -/* Finalize #pragma acc routine clauses after direct declarator has - been parsed, and put that into "omp declare target" attribute. */ - -static void -cp_parser_finish_oacc_routine (cp_parser *ARG_UNUSED (parser), tree fndecl, - tree clauses, bool named, bool is_defn, - bool first) -{ - location_t loc = OMP_CLAUSE_LOCATION (TREE_PURPOSE (clauses)); - - if (named && fndecl && is_overloaded_fn (fndecl) - && (TREE_CODE (fndecl) != FUNCTION_DECL - || DECL_FUNCTION_TEMPLATE_P (fndecl))) - { - error_at (loc, "%<#pragma acc routine%> names a set of overloads"); - return; - } - - if (!fndecl || TREE_CODE (fndecl) != FUNCTION_DECL - || (!named && !first)) - { - error_at (loc, "%<#pragma acc routine%> %s", - named ? "does not refer to a function" - : "not followed by single function"); - return; - } - - /* Perhaps we should use the same rule as declarations in different - namespaces? */ - if (named && !DECL_NAMESPACE_SCOPE_P (fndecl)) - { - error_at (loc, "%<#pragma acc routine%> does not refer to a" - " namespace scope function"); - return; - } - - if (get_oacc_fn_attrib (fndecl)) - error_at (loc, "%<#pragma acc routine%> already applied to %D", fndecl); - - if (TREE_USED (fndecl) || (!is_defn && DECL_SAVED_TREE (fndecl))) - error_at (OMP_CLAUSE_LOCATION (TREE_PURPOSE (clauses)), - "%<#pragma acc routine%> must be applied before %s", - TREE_USED (fndecl) ? "use" : "definition"); - - /* Process for function attrib */ - tree dims = build_oacc_routine_dims (TREE_VALUE (clauses)); - replace_oacc_fn_attrib (fndecl, dims); - - /* Also attach as a declare. */ - DECL_ATTRIBUTES (fndecl) - = tree_cons (get_identifier ("omp declare target"), - clauses, DECL_ATTRIBUTES (fndecl)); -} /* Parse the OpenACC routine pragma. This has an optional '( name )' component, which must resolve to a declared namespace-scope @@ -35837,16 +35783,50 @@ static void cp_parser_oacc_routine (cp_parser *parser, cp_token *pragma_tok, enum pragma_context context) { + bool first_p = parser->oacc_routine == NULL; + location_t loc = pragma_tok->location; + cp_omp_declare_simd_data data; + if (first_p) + { + data.error_seen = false; + data.fndecl_seen = false; + data.tokens = vNULL; + data.clauses = NULL_TREE; + parser->oacc_routine = &data; + } + tree decl = NULL_TREE; /* Create a dummy claue, to record location. */ tree c_head = build_omp_clause (pragma_tok->location, OMP_CLAUSE_SEQ); if (context != pragma_external) - cp_parser_error (parser, "%<#pragma acc routine%> not at file scope"); - + { + cp_parser_error (parser, "%<#pragma acc routine%> not at file scope"); + parser->oacc_routine->error_seen = true; + parser->oacc_routine = NULL; + return; + } + /* Look for optional '( name )'. */ - if (cp_lexer_next_token_is (parser->lexer,CPP_OPEN_PAREN)) + if (cp_lexer_next_token_is (parser->lexer, CPP_OPEN_PAREN)) { + if (!first_p) + { + while (cp_lexer_next_token_is_not (parser->lexer, CPP_PRAGMA_EOL) + && cp_lexer_next_token_is_not (parser->lexer, CPP_EOF)) + cp_lexer_consume_token (parser->lexer); + if (cp_lexer_next_token_is_not (parser->lexer, CPP_PRAGMA_EOL)) + parser->oacc_routine->error_seen = true; + cp_parser_require_pragma_eol (parser, pragma_tok); + + error_at (OMP_CLAUSE_LOCATION (parser->oacc_routine->clauses), + "%<#pragma oacc routine%> not followed by a single " + "function declaration or definition"); + + parser->oacc_routine->error_seen = true; + return; + } + cp_lexer_consume_token (parser->lexer); cp_token *token = cp_lexer_peek_token (parser->lexer); @@ -35868,36 +35848,192 @@ cp_parser_oacc_routine (cp_parser *parser, cp_token *pragma_tok, || !cp_parser_require (parser, CPP_CLOSE_PAREN, RT_CLOSE_PAREN)) { cp_parser_skip_to_pragma_eol (parser, pragma_tok); + parser->oacc_routine = NULL; return; } + + /* Build a chain of clauses. */ + parser->lexer->in_pragma = true; + tree clauses = NULL_TREE; + clauses = cp_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK, + "#pragma acc routine", + cp_lexer_peek_token + (parser->lexer)); + + /* Force clauses to be non-null, by attaching context to it. */ + clauses = tree_cons (c_head, clauses, NULL_TREE); + + if (decl && is_overloaded_fn (decl) + && (TREE_CODE (decl) != FUNCTION_DECL + || DECL_FUNCTION_TEMPLATE_P (decl))) + { + error_at (loc, "%<#pragma acc routine%> names a set of overloads"); + parser->oacc_routine = NULL; + return; + } + + /* Perhaps we should use the same rule as declarations in different + namespaces? */ + if (!DECL_NAMESPACE_SCOPE_P (decl)) + { + error_at (loc, "%<#pragma acc routine%> does not refer to a " + "namespace scope function"); + parser->oacc_routine = NULL; + return; + } + + if (!decl || TREE_CODE (decl) != FUNCTION_DECL) + { + error_at (loc, + "%<#pragma acc routine%> does not refer to a function"); + parser->oacc_routine = NULL; + return; + } + + data.clauses = clauses; + + cp_finalize_oacc_routine (parser, decl, false); + data.tokens.release (); + parser->oacc_routine = NULL; + } + else + { + while (cp_lexer_next_token_is_not (parser->lexer, CPP_PRAGMA_EOL) + && cp_lexer_next_token_is_not (parser->lexer, CPP_EOF)) + cp_lexer_consume_token (parser->lexer); + if (cp_lexer_next_token_is_not (parser->lexer, CPP_PRAGMA_EOL)) + parser->oacc_routine->error_seen = true; + cp_parser_require_pragma_eol (parser, pragma_tok); + + struct cp_token_cache *cp + = cp_token_cache_new (pragma_tok, cp_lexer_peek_token (parser->lexer)); + parser->oacc_routine->tokens.safe_push (cp); + + if (first_p) + parser->oacc_routine->clauses = c_head; + + while (cp_lexer_next_token_is (parser->lexer, CPP_PRAGMA)) + cp_parser_pragma (parser, context); + + if (first_p) + { + /* Create an empty list of clauses. */ + parser->oacc_routine->clauses = tree_cons (c_head, NULL_TREE, + NULL_TREE); + cp_parser_declaration (parser); + + if (parser->oacc_routine + && !parser->oacc_routine->error_seen + && !parser->oacc_routine->fndecl_seen) + error_at (loc, "%<#pragma acc routine%> not followed by " + "function declaration or definition"); + + data.tokens.release (); + parser->oacc_routine = NULL; + } + } +} + +/* Finalize #pragma acc routine clauses after direct declarator has + been parsed, and put that into "oacc routine" attribute. */ + +static tree +cp_parser_late_parsing_oacc_routine (cp_parser *parser, tree attrs) +{ + struct cp_token_cache *ce; + cp_omp_declare_simd_data *data = parser->oacc_routine; + tree cl, clauses = parser->oacc_routine->clauses; + location_t loc; + + loc = OMP_CLAUSE_LOCATION (TREE_PURPOSE(clauses)); + + if ((!data->error_seen && data->fndecl_seen) + || data->tokens.length () != 1) + { + error_at (loc, "%<#pragma oacc routine%> not followed by a single " + "function declaration or definition"); + data->error_seen = true; + return attrs; } + if (data->error_seen) + return attrs; + + ce = data->tokens[0]; - /* Build a chain of clauses. */ + cp_parser_push_lexer_for_tokens (parser, ce); parser->lexer->in_pragma = true; - tree clauses = NULL_TREE; - clauses = cp_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK, - "#pragma acc routine", - cp_lexer_peek_token (parser->lexer)); + gcc_assert (cp_lexer_peek_token (parser->lexer)->type == CPP_PRAGMA); + + cp_token *pragma_tok = cp_lexer_consume_token (parser->lexer); + cl = cp_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK, + "#pragma oacc routine", pragma_tok); + cp_parser_pop_lexer (parser); + + tree c_head = build_omp_clause (loc, OMP_CLAUSE_SEQ); /* Force clauses to be non-null, by attaching context to it. */ - clauses = tree_cons (c_head, clauses, NULL_TREE); + parser->oacc_routine->clauses = tree_cons (c_head, cl, NULL_TREE); - if (decl) - cp_parser_finish_oacc_routine (parser, decl, clauses, true, false, 0); - else - parser->oacc_routine = clauses; + data->fndecl_seen = true; + return attrs; } /* Apply any saved OpenACC routine clauses to a just-parsed declaration. */ static void -cp_finalize_oacc_routine (cp_parser *parser, tree fndecl, bool is_defn, - bool first) +cp_finalize_oacc_routine (cp_parser *parser, tree fndecl, bool is_defn) { - if (parser->oacc_routine) - cp_parser_finish_oacc_routine (parser, fndecl, parser->oacc_routine, - false, is_defn, first); + if (__builtin_expect (parser->oacc_routine != NULL, 0)) + { + tree clauses = parser->oacc_routine->clauses; + location_t loc = OMP_CLAUSE_LOCATION (TREE_PURPOSE(clauses)); + + if (parser->oacc_routine->error_seen) + return; + + if (fndecl == error_mark_node) + { + parser->oacc_routine = NULL; + return; + } + + if (TREE_CODE (fndecl) != FUNCTION_DECL) + { + cp_ensure_no_oacc_routine (parser); + return; + } + + if (!fndecl || TREE_CODE (fndecl) != FUNCTION_DECL) + { + error_at (loc, + "%<#pragma acc routine%> not followed by single function"); + parser->oacc_routine = NULL; + } + + if (get_oacc_fn_attrib (fndecl)) + { + error_at (loc, "%<#pragma acc routine%> already applied to %D", + fndecl); + parser->oacc_routine = NULL; + } + + if (TREE_USED (fndecl) || (!is_defn && DECL_SAVED_TREE (fndecl))) + { + error_at (loc, "%<#pragma acc routine%> must be applied before %s", + TREE_USED (fndecl) ? "use" : "definition"); + parser->oacc_routine = NULL; + } + + /* Process for function attrib */ + tree dims = build_oacc_routine_dims (TREE_VALUE (clauses)); + replace_oacc_fn_attrib (fndecl, dims); + + /* Add an "omp target" attribute. */ + DECL_ATTRIBUTES (fndecl) + = tree_cons (get_identifier ("omp declare target"), + NULL_TREE, DECL_ATTRIBUTES (fndecl)); + } } /* Main entry point to OpenMP statement pragmas. */ @@ -36381,7 +36517,6 @@ cp_parser_pragma (cp_parser *parser, enum pragma_context context) id = pragma_tok->pragma_kind; if (id != PRAGMA_OMP_DECLARE_REDUCTION && id != PRAGMA_OACC_ROUTINE) cp_ensure_no_omp_declare_simd (parser); - cp_ensure_no_oacc_routine (parser); switch (id) { case PRAGMA_GCC_PCH_PREPROCESS: diff --git a/gcc/cp/parser.h b/gcc/cp/parser.h index 022d037..a6b8e74 100644 --- a/gcc/cp/parser.h +++ b/gcc/cp/parser.h @@ -203,6 +203,7 @@ struct cp_omp_declare_simd_data { bool error_seen; /* Set if error has been reported. */ bool fndecl_seen; /* Set if one fn decl/definition has been seen already. */ vec<cp_token_cache_ptr> tokens; + tree clauses; }; @@ -371,8 +372,8 @@ struct GTY(()) cp_parser { necessary. */ cp_omp_declare_simd_data * GTY((skip)) cilk_simd_fn_info; - /* OpenACC routine clauses for subsequent decl/defn. */ - tree oacc_routine; + /* Parsing information for #pragma acc routine. */ + cp_omp_declare_simd_data * GTY((skip)) oacc_routine; /* Nonzero if parsing a parameter list where 'auto' should trigger an implicit template parameter. */ ^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [1/2] OpenACC routine support 2015-11-18 19:02 ` Cesar Philippidis @ 2015-11-19 11:23 ` Jakub Jelinek 2015-12-01 14:40 ` Thomas Schwinge 1 sibling, 0 replies; 23+ messages in thread From: Jakub Jelinek @ 2015-11-19 11:23 UTC (permalink / raw) To: Cesar Philippidis; +Cc: Nathan Sidwell, GCC Patches On Wed, Nov 18, 2015 at 11:02:01AM -0800, Cesar Philippidis wrote: > static inline void > cp_ensure_no_oacc_routine (cp_parser *parser) > { > - cp_finalize_oacc_routine (parser, NULL_TREE, false, true); > + if (parser->oacc_routine && !parser->oacc_routine->error_seen) > + { > + tree clauses = parser->oacc_routine->clauses; > + location_t loc = OMP_CLAUSE_LOCATION (TREE_PURPOSE(clauses)); Formatting, missing space before (clauses)); > @@ -19326,7 +19330,11 @@ cp_parser_late_return_type_opt (cp_parser* parser, cp_declarator *declarator, > declarator->std_attributes > = cp_parser_late_parsing_omp_declare_simd (parser, > declarator->std_attributes); > - > + if (oacc_routine_p) > + declarator->std_attributes > + = cp_parser_late_parsing_oacc_routine (parser, > + declarator->std_attributes); > + Trailing whitespace at the end of line. Otherwise LGTM. Jakub ^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [1/2] OpenACC routine support 2015-11-18 19:02 ` Cesar Philippidis 2015-11-19 11:23 ` Jakub Jelinek @ 2015-12-01 14:40 ` Thomas Schwinge 2015-12-01 14:49 ` Cesar Philippidis 1 sibling, 1 reply; 23+ messages in thread From: Thomas Schwinge @ 2015-12-01 14:40 UTC (permalink / raw) To: Cesar Philippidis; +Cc: Nathan Sidwell, GCC Patches, Jakub Jelinek [-- Attachment #1: Type: text/plain, Size: 2714 bytes --] Hi Cesar! I noticed while working on other test cases: On Wed, 18 Nov 2015 11:02:01 -0800, Cesar Philippidis <cesar@codesourcery.com> wrote: > --- a/gcc/cp/parser.c > +++ b/gcc/cp/parser.c > @@ -1318,13 +1318,21 @@ cp_finalize_omp_declare_simd (cp_parser *parser, tree fndecl) > } > } > > -/* Diagnose if #pragma omp routine isn't followed immediately > - by function declaration or definition. */ > +/* Diagnose if #pragma acc routine isn't followed immediately by function > + declaration or definition. */ > > static inline void > cp_ensure_no_oacc_routine (cp_parser *parser) > { > - cp_finalize_oacc_routine (parser, NULL_TREE, false, true); > + if (parser->oacc_routine && !parser->oacc_routine->error_seen) > + { > + tree clauses = parser->oacc_routine->clauses; > + location_t loc = OMP_CLAUSE_LOCATION (TREE_PURPOSE(clauses)); > + > + error_at (loc, "%<#pragma oacc routine%> not followed by function " > + "declaration or definition"); > + parser->oacc_routine = NULL; > + } > } "#pragma acc routine", not "oacc". Also in a few other places. Next, in the function quoted above, you use "not followed by function declaration or definition", but you use "not followed by a single function declaration or definition" in a lot of (but not all) other places -- is that intentional? For example: > cp_parser_oacc_routine (cp_parser *parser, cp_token *pragma_tok, > enum pragma_context context) > { > [...] > + error_at (OMP_CLAUSE_LOCATION (parser->oacc_routine->clauses), > + "%<#pragma oacc routine%> not followed by a single " > + "function declaration or definition"); "a single". > [...] > + if (parser->oacc_routine > + && !parser->oacc_routine->error_seen > + && !parser->oacc_routine->fndecl_seen) > + error_at (loc, "%<#pragma acc routine%> not followed by " > + "function declaration or definition"); Not "a single". > + > + data.tokens.release (); > + parser->oacc_routine = NULL; > + } > + } > +} > + > +/* Finalize #pragma acc routine clauses after direct declarator has > + been parsed, and put that into "oacc routine" attribute. */ There is no "oacc routine" attribute (anymore)? > +static tree > +cp_parser_late_parsing_oacc_routine (cp_parser *parser, tree attrs) > +{ > [...] > + if ((!data->error_seen && data->fndecl_seen) > + || data->tokens.length () != 1) > + { > + error_at (loc, "%<#pragma oacc routine%> not followed by a single " > + "function declaration or definition"); "a single". (I have not verified all of the parser(s) source code.) Grüße Thomas [-- Attachment #2: signature.asc --] [-- Type: application/pgp-signature, Size: 472 bytes --] ^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [1/2] OpenACC routine support 2015-12-01 14:40 ` Thomas Schwinge @ 2015-12-01 14:49 ` Cesar Philippidis 2015-12-02 23:37 ` Cesar Philippidis 0 siblings, 1 reply; 23+ messages in thread From: Cesar Philippidis @ 2015-12-01 14:49 UTC (permalink / raw) To: Thomas Schwinge; +Cc: Nathan Sidwell, GCC Patches, Jakub Jelinek On 12/01/2015 06:40 AM, Thomas Schwinge wrote: > I noticed while working on other test cases: > > On Wed, 18 Nov 2015 11:02:01 -0800, Cesar Philippidis <cesar@codesourcery.com> wrote: >> --- a/gcc/cp/parser.c >> +++ b/gcc/cp/parser.c > >> @@ -1318,13 +1318,21 @@ cp_finalize_omp_declare_simd (cp_parser *parser, tree fndecl) >> } >> } >> >> -/* Diagnose if #pragma omp routine isn't followed immediately >> - by function declaration or definition. */ >> +/* Diagnose if #pragma acc routine isn't followed immediately by function >> + declaration or definition. */ >> >> static inline void >> cp_ensure_no_oacc_routine (cp_parser *parser) >> { >> - cp_finalize_oacc_routine (parser, NULL_TREE, false, true); >> + if (parser->oacc_routine && !parser->oacc_routine->error_seen) >> + { >> + tree clauses = parser->oacc_routine->clauses; >> + location_t loc = OMP_CLAUSE_LOCATION (TREE_PURPOSE(clauses)); >> + >> + error_at (loc, "%<#pragma oacc routine%> not followed by function " >> + "declaration or definition"); >> + parser->oacc_routine = NULL; >> + } >> } > > "#pragma acc routine", not "oacc". Also in a few other places. Good eyes. Thanks for catching that. > Next, in the function quoted above, you use "not followed by function > declaration or definition", but you use "not followed by a single > function declaration or definition" in a lot of (but not all) other > places -- is that intentional? I probably wasn't being consistent. Which error message do you prefer? I'll take a look at what the c front end does. > For example: > >> cp_parser_oacc_routine (cp_parser *parser, cp_token *pragma_tok, >> enum pragma_context context) >> { >> [...] >> + error_at (OMP_CLAUSE_LOCATION (parser->oacc_routine->clauses), >> + "%<#pragma oacc routine%> not followed by a single " >> + "function declaration or definition"); > > "a single". > >> [...] >> + if (parser->oacc_routine >> + && !parser->oacc_routine->error_seen >> + && !parser->oacc_routine->fndecl_seen) >> + error_at (loc, "%<#pragma acc routine%> not followed by " >> + "function declaration or definition"); > > Not "a single". > >> + >> + data.tokens.release (); >> + parser->oacc_routine = NULL; >> + } >> + } >> +} >> + >> +/* Finalize #pragma acc routine clauses after direct declarator has >> + been parsed, and put that into "oacc routine" attribute. */ > > There is no "oacc routine" attribute (anymore)? You're right, it was renamed to 'oacc function'. >> +static tree >> +cp_parser_late_parsing_oacc_routine (cp_parser *parser, tree attrs) >> +{ >> [...] >> + if ((!data->error_seen && data->fndecl_seen) >> + || data->tokens.length () != 1) >> + { >> + error_at (loc, "%<#pragma oacc routine%> not followed by a single " >> + "function declaration or definition"); > > "a single". > > (I have not verified all of the parser(s) source code.) Thanks. I'll go through and update the comments and error messages. Cesar ^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [1/2] OpenACC routine support 2015-12-01 14:49 ` Cesar Philippidis @ 2015-12-02 23:37 ` Cesar Philippidis 2015-12-03 8:36 ` Thomas Schwinge 0 siblings, 1 reply; 23+ messages in thread From: Cesar Philippidis @ 2015-12-02 23:37 UTC (permalink / raw) To: Thomas Schwinge; +Cc: Nathan Sidwell, GCC Patches, Jakub Jelinek On 12/01/2015 06:49 AM, Cesar Philippidis wrote: > On 12/01/2015 06:40 AM, Thomas Schwinge wrote: > >> I noticed while working on other test cases: >> >> On Wed, 18 Nov 2015 11:02:01 -0800, Cesar Philippidis <cesar@codesourcery.com> wrote: >>> --- a/gcc/cp/parser.c >>> +++ b/gcc/cp/parser.c >> >>> @@ -1318,13 +1318,21 @@ cp_finalize_omp_declare_simd (cp_parser *parser, tree fndecl) >>> } >>> } >>> >>> -/* Diagnose if #pragma omp routine isn't followed immediately >>> - by function declaration or definition. */ >>> +/* Diagnose if #pragma acc routine isn't followed immediately by function >>> + declaration or definition. */ >>> >>> static inline void >>> cp_ensure_no_oacc_routine (cp_parser *parser) >>> { >>> - cp_finalize_oacc_routine (parser, NULL_TREE, false, true); >>> + if (parser->oacc_routine && !parser->oacc_routine->error_seen) >>> + { >>> + tree clauses = parser->oacc_routine->clauses; >>> + location_t loc = OMP_CLAUSE_LOCATION (TREE_PURPOSE(clauses)); >>> + >>> + error_at (loc, "%<#pragma oacc routine%> not followed by function " >>> + "declaration or definition"); >>> + parser->oacc_routine = NULL; >>> + } >>> } >> >> "#pragma acc routine", not "oacc". Also in a few other places. > > Good eyes. Thanks for catching that. > >> Next, in the function quoted above, you use "not followed by function >> declaration or definition", but you use "not followed by a single >> function declaration or definition" in a lot of (but not all) other >> places -- is that intentional? > > I probably wasn't being consistent. Which error message do you prefer? > I'll take a look at what the c front end does. > >> For example: >> >>> cp_parser_oacc_routine (cp_parser *parser, cp_token *pragma_tok, >>> enum pragma_context context) >>> { >>> [...] >>> + error_at (OMP_CLAUSE_LOCATION (parser->oacc_routine->clauses), >>> + "%<#pragma oacc routine%> not followed by a single " >>> + "function declaration or definition"); >> >> "a single". >> >>> [...] >>> + if (parser->oacc_routine >>> + && !parser->oacc_routine->error_seen >>> + && !parser->oacc_routine->fndecl_seen) >>> + error_at (loc, "%<#pragma acc routine%> not followed by " >>> + "function declaration or definition"); >> >> Not "a single". >> >>> + >>> + data.tokens.release (); >>> + parser->oacc_routine = NULL; >>> + } >>> + } >>> +} >>> + >>> +/* Finalize #pragma acc routine clauses after direct declarator has >>> + been parsed, and put that into "oacc routine" attribute. */ >> >> There is no "oacc routine" attribute (anymore)? > > You're right, it was renamed to 'oacc function'. > >>> +static tree >>> +cp_parser_late_parsing_oacc_routine (cp_parser *parser, tree attrs) >>> +{ >>> [...] >>> + if ((!data->error_seen && data->fndecl_seen) >>> + || data->tokens.length () != 1) >>> + { >>> + error_at (loc, "%<#pragma oacc routine%> not followed by a single " >>> + "function declaration or definition"); >> >> "a single". >> >> (I have not verified all of the parser(s) source code.) > > Thanks. I'll go through and update the comments and error messages. Here's the updated patch. The test cases were written in a way such that none of them needed to be updated with these changes. I'm tempted to commit this as obvious, but I want to make sure you're ok with these new messages. The major change is to report these errors as "pragma acc routine not followed by a function declaration or definition". I think that's more descriptive then "not followed by a single function". That said, it looks like the c front end uses the latter error message. Is this OK or do you prefer the "not followed by a single function" message? Cesar ^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [1/2] OpenACC routine support 2015-12-02 23:37 ` Cesar Philippidis @ 2015-12-03 8:36 ` Thomas Schwinge 2015-12-03 14:06 ` Cesar Philippidis 0 siblings, 1 reply; 23+ messages in thread From: Thomas Schwinge @ 2015-12-03 8:36 UTC (permalink / raw) To: Cesar Philippidis; +Cc: Nathan Sidwell, GCC Patches, Jakub Jelinek [-- Attachment #1: Type: text/plain, Size: 3075 bytes --] Hi Cesar! On Wed, 2 Dec 2015 15:37:17 -0800, Cesar Philippidis <cesar_philippidis@mentor.com> wrote: > On 12/01/2015 06:49 AM, Cesar Philippidis wrote: > > On 12/01/2015 06:40 AM, Thomas Schwinge wrote: > > > >> I noticed while working on other test cases: > >> > >> On Wed, 18 Nov 2015 11:02:01 -0800, Cesar Philippidis <cesar@codesourcery.com> wrote: > >>> --- a/gcc/cp/parser.c > >>> +++ b/gcc/cp/parser.c > >> > >>> @@ -1318,13 +1318,21 @@ cp_finalize_omp_declare_simd (cp_parser *parser, tree fndecl) > >>> } > >>> } > >>> > >>> -/* Diagnose if #pragma omp routine isn't followed immediately > >>> - by function declaration or definition. */ > >>> +/* Diagnose if #pragma acc routine isn't followed immediately by function > >>> + declaration or definition. */ > >>> > >>> static inline void > >>> cp_ensure_no_oacc_routine (cp_parser *parser) > >>> { > >>> - cp_finalize_oacc_routine (parser, NULL_TREE, false, true); > >>> + if (parser->oacc_routine && !parser->oacc_routine->error_seen) > >>> + { > >>> + tree clauses = parser->oacc_routine->clauses; > >>> + location_t loc = OMP_CLAUSE_LOCATION (TREE_PURPOSE(clauses)); > >>> + > >>> + error_at (loc, "%<#pragma oacc routine%> not followed by function " > >>> + "declaration or definition"); > >>> + parser->oacc_routine = NULL; > >>> + } > >>> } > >> Next, in the function quoted above, you use "not followed by function > >> declaration or definition", but you use "not followed by a single > >> function declaration or definition" in a lot of (but not all) other > >> places -- is that intentional? > > > > I probably wasn't being consistent. Which error message do you prefer? > > I'll take a look at what the c front end does. > > > >> For example: [...] > >> (I have not verified all of the parser(s) source code.) > > > > Thanks. I'll go through and update the comments and error messages. > > Here's the updated patch. ENOPATCH. > The test cases were written in a way such that > none of them needed to be updated with these changes. ... which potentially means they'd match for all kinds of "random" diagnostics. ;-) > I'm tempted to commit this as obvious, but I want to make sure you're ok > with these new messages. I don't care very much, as long as it's understandable for a user. I just tripped over this because of mismatches between C and C++ as well as different C++ diagnostic variants. > The major change is to report these errors as > "pragma acc routine not followed by a function declaration or > definition". I think that's more descriptive then "not followed by a > single function". That said, it looks like the c front end uses the > latter error message. (In the C front end, the "a" is missing: "not followed by single function"; that should be fixed up as well.) > Is this OK or do you prefer the "not followed by a single function" message? "not followed by a function declaration or definition" sounds good to me. Grüße Thomas [-- Attachment #2: signature.asc --] [-- Type: application/pgp-signature, Size: 472 bytes --] ^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [1/2] OpenACC routine support 2015-12-03 8:36 ` Thomas Schwinge @ 2015-12-03 14:06 ` Cesar Philippidis 0 siblings, 0 replies; 23+ messages in thread From: Cesar Philippidis @ 2015-12-03 14:06 UTC (permalink / raw) To: Thomas Schwinge; +Cc: Nathan Sidwell, GCC Patches, Jakub Jelinek [-- Attachment #1: Type: text/plain, Size: 1334 bytes --] On 12/03/2015 12:36 AM, Thomas Schwinge wrote: >> Here's the updated patch. > > ENOPATCH. Here it is. >> The test cases were written in a way such that >> none of them needed to be updated with these changes. > > ... which potentially means they'd match for all kinds of "random" > diagnostics. ;-) They were supposed to be generic enough so that they work both in c and c++. But, yeah, that randomness is likely. >> I'm tempted to commit this as obvious, but I want to make sure you're ok >> with these new messages. > > I don't care very much, as long as it's understandable for a user. I > just tripped over this because of mismatches between C and C++ as well as > different C++ diagnostic variants. > >> The major change is to report these errors as >> "pragma acc routine not followed by a function declaration or >> definition". I think that's more descriptive then "not followed by a >> single function". That said, it looks like the c front end uses the >> latter error message. > > (In the C front end, the "a" is missing: "not followed by single > function"; that should be fixed up as well.) > >> Is this OK or do you prefer the "not followed by a single function" message? > > "not followed by a function declaration or definition" sounds good to me. Ok, I'll apply this patch in a couple of hours. Cesar [-- Attachment #2: cxx-routines-cleanup.diff --] [-- Type: text/x-patch, Size: 3283 bytes --] 2015-12-02 Cesar Philippidis <cesar@codesourcery.com> gcc/cp/ * parser.c (cp_ensure_no_oacc_routine): Update error message. (cp_parser_oacc_routine): Likewise. (cp_parser_late_parsing_oacc_routine): Likewise. Update comment describing this function. (cp_finalize_oacc_routine): Update error message. diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index b4ecac7..1c14354 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -1329,7 +1329,7 @@ cp_ensure_no_oacc_routine (cp_parser *parser) tree clauses = parser->oacc_routine->clauses; location_t loc = OMP_CLAUSE_LOCATION (TREE_PURPOSE (clauses)); - error_at (loc, "%<#pragma oacc routine%> not followed by function " + error_at (loc, "%<#pragma acc routine%> not followed by a function " "declaration or definition"); parser->oacc_routine = NULL; } @@ -35857,7 +35857,7 @@ cp_parser_oacc_routine (cp_parser *parser, cp_token *pragma_tok, cp_parser_require_pragma_eol (parser, pragma_tok); error_at (OMP_CLAUSE_LOCATION (parser->oacc_routine->clauses), - "%<#pragma oacc routine%> not followed by a single " + "%<#pragma acc routine%> not followed by a " "function declaration or definition"); parser->oacc_routine->error_seen = true; @@ -35962,7 +35962,7 @@ cp_parser_oacc_routine (cp_parser *parser, cp_token *pragma_tok, if (parser->oacc_routine && !parser->oacc_routine->error_seen && !parser->oacc_routine->fndecl_seen) - error_at (loc, "%<#pragma acc routine%> not followed by " + error_at (loc, "%<#pragma acc routine%> not followed by a " "function declaration or definition"); data.tokens.release (); @@ -35972,7 +35972,7 @@ cp_parser_oacc_routine (cp_parser *parser, cp_token *pragma_tok, } /* Finalize #pragma acc routine clauses after direct declarator has - been parsed, and put that into "oacc routine" attribute. */ + been parsed, and put that into "oacc function" attribute. */ static tree cp_parser_late_parsing_oacc_routine (cp_parser *parser, tree attrs) @@ -35987,7 +35987,7 @@ cp_parser_late_parsing_oacc_routine (cp_parser *parser, tree attrs) if ((!data->error_seen && data->fndecl_seen) || data->tokens.length () != 1) { - error_at (loc, "%<#pragma oacc routine%> not followed by a single " + error_at (loc, "%<#pragma acc routine%> not followed by a " "function declaration or definition"); data->error_seen = true; return attrs; @@ -36003,7 +36003,7 @@ cp_parser_late_parsing_oacc_routine (cp_parser *parser, tree attrs) cp_token *pragma_tok = cp_lexer_consume_token (parser->lexer); cl = cp_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK, - "#pragma oacc routine", pragma_tok); + "#pragma acc routine", pragma_tok); cp_parser_pop_lexer (parser); tree c_head = build_omp_clause (loc, OMP_CLAUSE_SEQ); @@ -36044,7 +36044,8 @@ cp_finalize_oacc_routine (cp_parser *parser, tree fndecl, bool is_defn) if (!fndecl || TREE_CODE (fndecl) != FUNCTION_DECL) { error_at (loc, - "%<#pragma acc routine%> not followed by single function"); + "%<#pragma acc routine%> not followed by a function " + "declaration or definition"); parser->oacc_routine = NULL; } ^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [2/2] OpenACC routine support 2015-11-02 18:56 [0/2] OpenACC routine support Nathan Sidwell 2015-11-02 19:21 ` [1/2] " Nathan Sidwell @ 2015-11-02 19:23 ` Nathan Sidwell 2015-11-02 19:41 ` Jakub Jelinek 2015-11-03 15:38 ` Jakub Jelinek 1 sibling, 2 replies; 23+ messages in thread From: Nathan Sidwell @ 2015-11-02 19:23 UTC (permalink / raw) To: Jakub Jelinek, GCC Patches [-- Attachment #1: Type: text/plain, Size: 308 bytes --] Here are the tests for the routine support. The compiler tests check invalid combinations of gang, worker, vector & seq. The libgomp execution tests check the expected partioning occurs within loops. As with the reduction tests, these ones are taken from the execution model loop tests. ok? nathan [-- Attachment #2: 02-trunk-routine-tests-1102.patch --] [-- Type: text/x-patch, Size: 12267 bytes --] 2015-11-02 Nathan Sidwell <nathan@codesourcery.com> gcc/testsuite/ * c-c++-common/goacc/routine-1.c: New. * c-c++-common/goacc/routine-2.c: New. * c-c++-common/goacc/routine-3.c: New. * c-c++-common/goacc/routine-4.c: New. libgomp/ * testsuite/libgomp.oacc-c-c++-common/routine-g-1.c: New. * testsuite/libgomp.oacc-c-c++-common/routine-gwv-1.c: New. * testsuite/libgomp.oacc-c-c++-common/routine-v-1.c: New. * testsuite/libgomp.oacc-c-c++-common/routine-w-1.c: New. * testsuite/libgomp.oacc-c-c++-common/routine-wv-1.c: New. Index: gcc/testsuite/c-c++-common/goacc/routine-1.c =================================================================== --- gcc/testsuite/c-c++-common/goacc/routine-1.c (revision 0) +++ gcc/testsuite/c-c++-common/goacc/routine-1.c (working copy) @@ -0,0 +1,34 @@ + +#pragma acc routine gang +void gang (void) +{ +} + +#pragma acc routine worker +void worker (void) +{ +} + +#pragma acc routine vector +void vector (void) +{ +} + +#pragma acc routine seq +void seq (void) +{ +} + +int main () +{ + +#pragma acc parallel num_gangs (32) num_workers (32) vector_length (32) + { + gang (); + worker (); + vector (); + seq (); + } + + return 0; +} Index: gcc/testsuite/c-c++-common/goacc/routine-2.c =================================================================== --- gcc/testsuite/c-c++-common/goacc/routine-2.c (revision 0) +++ gcc/testsuite/c-c++-common/goacc/routine-2.c (working copy) @@ -0,0 +1,21 @@ +#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) +{ +} + +#pragma acc routine (nothing) gang /* { dg-error "not been declared" } */ Index: gcc/testsuite/c-c++-common/goacc/routine-3.c =================================================================== --- gcc/testsuite/c-c++-common/goacc/routine-3.c (revision 0) +++ gcc/testsuite/c-c++-common/goacc/routine-3.c (working copy) @@ -0,0 +1,53 @@ +#pragma acc routine gang +void gang (void) /* { dg-message "declared here" 3 } */ +{ +} + +#pragma acc routine worker +void worker (void) /* { dg-message "declared here" 2 } */ +{ +} + +#pragma acc routine vector +void vector (void) /* { dg-message "declared here" 1 } */ +{ +} + +#pragma acc routine seq +void seq (void) +{ +} + +int main () +{ + +#pragma acc parallel num_gangs (32) num_workers (32) vector_length (32) + { + #pragma acc loop gang /* { dg-message "loop here" 1 } */ + for (int i = 0; i < 10; i++) + { + gang (); /* { dg-error "routine call uses same" } */ + worker (); + vector (); + seq (); + } + #pragma acc loop worker /* { dg-message "loop here" 2 } */ + for (int i = 0; i < 10; i++) + { + gang (); /* { dg-error "routine call uses same" } */ + worker (); /* { dg-error "routine call uses same" } */ + vector (); + seq (); + } + #pragma acc loop vector /* { dg-message "loop here" 3 } */ + for (int i = 0; i < 10; i++) + { + gang (); /* { dg-error "routine call uses same" } */ + worker (); /* { dg-error "routine call uses same" } */ + vector (); /* { dg-error "routine call uses same" } */ + seq (); + } + } + + return 0; +} Index: gcc/testsuite/c-c++-common/goacc/routine-4.c =================================================================== --- gcc/testsuite/c-c++-common/goacc/routine-4.c (revision 0) +++ gcc/testsuite/c-c++-common/goacc/routine-4.c (working copy) @@ -0,0 +1,41 @@ + +void gang (void); +void worker (void); +void vector (void); + +#pragma acc routine (gang) gang +#pragma acc routine (worker) worker +#pragma acc routine (vector) vector + +#pragma acc routine seq +void seq (void) +{ + gang (); /* { dg-error "routine call uses" } */ + worker (); /* { dg-error "routine call uses" } */ + vector (); /* { dg-error "routine call uses" } */ + seq (); +} + +void vector (void) /* { dg-message "declared here" 1 } */ +{ + gang (); /* { dg-error "routine call uses" } */ + worker (); /* { dg-error "routine call uses" } */ + vector (); + seq (); +} + +void worker (void) /* { dg-message "declared here" 2 } */ +{ + gang (); /* { dg-error "routine call uses" } */ + worker (); + vector (); + seq (); +} + +void gang (void) /* { dg-message "declared here" 3 } */ +{ + gang (); + worker (); + vector (); + seq (); +} Index: libgomp/testsuite/libgomp.oacc-c-c++-common/routine-g-1.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/routine-g-1.c (revision 0) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/routine-g-1.c (working copy) @@ -0,0 +1,64 @@ +/* { dg-do run } */ +/* { dg-additional-options "-O2" */ + +#include <stdio.h> + +#define N (32*32*32+17) + +#pragma acc routine gang +void __attribute__ ((noinline)) gang (int ary[N]) +{ +#pragma acc loop gang + for (unsigned ix = 0; ix < N; ix++) + { + if (__builtin_acc_on_device (5)) + { + int g = 0, w = 0, v = 0; + + __asm__ volatile ("mov.u32 %0,%%ctaid.x;" : "=r" (g)); + __asm__ volatile ("mov.u32 %0,%%tid.y;" : "=r" (w)); + __asm__ volatile ("mov.u32 %0,%%tid.x;" : "=r" (v)); + ary[ix] = (g << 16) | (w << 8) | v; + } + else + ary[ix] = ix; + } +} + +int main () +{ + int ary[N]; + int ix; + int exit = 0; + int ondev = 0; + + for (ix = 0; ix < N;ix++) + ary[ix] = -1; + +#pragma acc parallel num_gangs(32) vector_length(32) copy(ary) copy(ondev) + { + ondev = __builtin_acc_on_device (5); + gang (ary); + } + + for (ix = 0; ix < N; ix++) + { + int expected = ix; + if(ondev) + { + int g = ix / ((N + 31) / 32); + int w = 0; + int v = 0; + + expected = (g << 16) | (w << 8) | v; + } + + if (ary[ix] != expected) + { + exit = 1; + printf ("ary[%d]=%x expected %x\n", ix, ary[ix], expected); + } + } + + return exit; +} Index: libgomp/testsuite/libgomp.oacc-c-c++-common/routine-gwv-1.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/routine-gwv-1.c (revision 0) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/routine-gwv-1.c (working copy) @@ -0,0 +1,66 @@ +/* { dg-do run } */ +/* { dg-additional-options "-O2" */ + +#include <stdio.h> + +#define N (32*32*32+17) + +#pragma acc routine gang +void __attribute__ ((noinline)) gang (int ary[N]) +{ +#pragma acc loop gang worker vector + for (unsigned ix = 0; ix < N; ix++) + { + if (__builtin_acc_on_device (5)) + { + int g = 0, w = 0, v = 0; + + __asm__ volatile ("mov.u32 %0,%%ctaid.x;" : "=r" (g)); + __asm__ volatile ("mov.u32 %0,%%tid.y;" : "=r" (w)); + __asm__ volatile ("mov.u32 %0,%%tid.x;" : "=r" (v)); + ary[ix] = (g << 16) | (w << 8) | v; + } + else + ary[ix] = ix; + } +} + +int main () +{ + int ary[N]; + int ix; + int exit = 0; + int ondev = 0; + + for (ix = 0; ix < N;ix++) + ary[ix] = -1; + +#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) copy(ary) copy(ondev) + { + ondev = __builtin_acc_on_device (5); + gang (ary); + } + + for (ix = 0; ix < N; ix++) + { + int expected = ix; + if(ondev) + { + int chunk_size = (N + 32*32*32 - 1) / (32*32*32); + + int g = ix / (chunk_size * 32 * 32); + int w = ix / 32 % 32; + int v = ix % 32; + + expected = (g << 16) | (w << 8) | v; + } + + if (ary[ix] != expected) + { + exit = 1; + printf ("ary[%d]=%x expected %x\n", ix, ary[ix], expected); + } + } + + return exit; +} Index: libgomp/testsuite/libgomp.oacc-c-c++-common/routine-v-1.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/routine-v-1.c (revision 0) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/routine-v-1.c (working copy) @@ -0,0 +1,64 @@ +/* { dg-do run } */ +/* { dg-additional-options "-O2" */ + +#include <stdio.h> + +#define N (32*32*32+17) + +#pragma acc routine vector +void __attribute__ ((noinline)) vector (int ary[N]) +{ +#pragma acc loop vector + for (unsigned ix = 0; ix < N; ix++) + { + if (__builtin_acc_on_device (5)) + { + int g = 0, w = 0, v = 0; + + __asm__ volatile ("mov.u32 %0,%%ctaid.x;" : "=r" (g)); + __asm__ volatile ("mov.u32 %0,%%tid.y;" : "=r" (w)); + __asm__ volatile ("mov.u32 %0,%%tid.x;" : "=r" (v)); + ary[ix] = (g << 16) | (w << 8) | v; + } + else + ary[ix] = ix; + } +} + +int main () +{ + int ary[N]; + int ix; + int exit = 0; + int ondev = 0; + + for (ix = 0; ix < N;ix++) + ary[ix] = -1; + +#pragma acc parallel vector_length(32) copy(ary) copy(ondev) + { + ondev = __builtin_acc_on_device (5); + vector (ary); + } + + for (ix = 0; ix < N; ix++) + { + int expected = ix; + if(ondev) + { + int g = 0; + int w = 0; + int v = ix % 32; + + expected = (g << 16) | (w << 8) | v; + } + + if (ary[ix] != expected) + { + exit = 1; + printf ("ary[%d]=%x expected %x\n", ix, ary[ix], expected); + } + } + + return exit; +} Index: libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c (revision 0) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c (working copy) @@ -0,0 +1,64 @@ +/* { dg-do run } */ +/* { dg-additional-options "-O2" */ + +#include <stdio.h> + +#define N (32*32*32+17) + +#pragma acc routine worker +void __attribute__ ((noinline)) worker (int ary[N]) +{ +#pragma acc loop worker + for (unsigned ix = 0; ix < N; ix++) + { + if (__builtin_acc_on_device (5)) + { + int g = 0, w = 0, v = 0; + + __asm__ volatile ("mov.u32 %0,%%ctaid.x;" : "=r" (g)); + __asm__ volatile ("mov.u32 %0,%%tid.y;" : "=r" (w)); + __asm__ volatile ("mov.u32 %0,%%tid.x;" : "=r" (v)); + ary[ix] = (g << 16) | (w << 8) | v; + } + else + ary[ix] = ix; + } +} + +int main () +{ + int ary[N]; + int ix; + int exit = 0; + int ondev = 0; + + for (ix = 0; ix < N;ix++) + ary[ix] = -1; + +#pragma acc parallel num_workers(32) vector_length(32) copy(ary) copy(ondev) + { + ondev = __builtin_acc_on_device (5); + worker (ary); + } + + for (ix = 0; ix < N; ix++) + { + int expected = ix; + if(ondev) + { + int g = 0; + int w = ix % 32; + int v = 0; + + expected = (g << 16) | (w << 8) | v; + } + + if (ary[ix] != expected) + { + exit = 1; + printf ("ary[%d]=%x expected %x\n", ix, ary[ix], expected); + } + } + + return exit; +} Index: libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-1.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-1.c (revision 0) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-1.c (working copy) @@ -0,0 +1,64 @@ +/* { dg-do run } */ +/* { dg-additional-options "-O2" */ + +#include <stdio.h> + +#define N (32*32*32+17) + +#pragma acc routine worker +void __attribute__ ((noinline)) worker (int ary[N]) +{ +#pragma acc loop worker vector + for (unsigned ix = 0; ix < N; ix++) + { + if (__builtin_acc_on_device (5)) + { + int g = 0, w = 0, v = 0; + + __asm__ volatile ("mov.u32 %0,%%ctaid.x;" : "=r" (g)); + __asm__ volatile ("mov.u32 %0,%%tid.y;" : "=r" (w)); + __asm__ volatile ("mov.u32 %0,%%tid.x;" : "=r" (v)); + ary[ix] = (g << 16) | (w << 8) | v; + } + else + ary[ix] = ix; + } +} + +int main () +{ + int ary[N]; + int ix; + int exit = 0; + int ondev = 0; + + for (ix = 0; ix < N;ix++) + ary[ix] = -1; + +#pragma acc parallel num_workers(32) vector_length(32) copy(ary) copy(ondev) + { + ondev = __builtin_acc_on_device (5); + worker (ary); + } + + for (ix = 0; ix < N; ix++) + { + int expected = ix; + if(ondev) + { + int g = 0; + int w = (ix / 32) % 32; + int v = ix % 32; + + expected = (g << 16) | (w << 8) | v; + } + + if (ary[ix] != expected) + { + exit = 1; + printf ("ary[%d]=%x expected %x\n", ix, ary[ix], expected); + } + } + + return exit; +} ^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [2/2] OpenACC routine support 2015-11-02 19:23 ` [2/2] " Nathan Sidwell @ 2015-11-02 19:41 ` Jakub Jelinek 2015-11-02 20:01 ` Nathan Sidwell 2015-11-03 15:38 ` Jakub Jelinek 1 sibling, 1 reply; 23+ messages in thread From: Jakub Jelinek @ 2015-11-02 19:41 UTC (permalink / raw) To: Nathan Sidwell; +Cc: GCC Patches On Mon, Nov 02, 2015 at 02:23:19PM -0500, Nathan Sidwell wrote: > +#pragma acc routine gang > +void __attribute__ ((noinline)) gang (int ary[N]) > +{ > +#pragma acc loop gang > + for (unsigned ix = 0; ix < N; ix++) > + { > + if (__builtin_acc_on_device (5)) > + { > + int g = 0, w = 0, v = 0; > + > + __asm__ volatile ("mov.u32 %0,%%ctaid.x;" : "=r" (g)); > + __asm__ volatile ("mov.u32 %0,%%tid.y;" : "=r" (w)); > + __asm__ volatile ("mov.u32 %0,%%tid.x;" : "=r" (v)); > + ary[ix] = (g << 16) | (w << 8) | v; > + } > + else > + ary[ix] = ix; Does this work even with -O0? I mean, the assembler is invalid for any target other than PTX, so you are relying on aggressively folding this away. Jakub ^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [2/2] OpenACC routine support 2015-11-02 19:41 ` Jakub Jelinek @ 2015-11-02 20:01 ` Nathan Sidwell 0 siblings, 0 replies; 23+ messages in thread From: Nathan Sidwell @ 2015-11-02 20:01 UTC (permalink / raw) To: Jakub Jelinek; +Cc: GCC Patches On 11/02/15 14:41, Jakub Jelinek wrote: > > Does this work even with -O0? I mean, the assembler is invalid > for any target other than PTX, so you are relying on aggressively folding > this away. Correct. As thread identification is inherently target-specific, I don't see how to do otherwise. We know _builtin_acc_on_device is folded for a constant arg, and -O2 enables dead code elimination such that non-PTX targets (such as the host) don't see that assembler. nathan ^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [2/2] OpenACC routine support 2015-11-02 19:23 ` [2/2] " Nathan Sidwell 2015-11-02 19:41 ` Jakub Jelinek @ 2015-11-03 15:38 ` Jakub Jelinek 2015-11-03 15:56 ` Nathan Sidwell 1 sibling, 1 reply; 23+ messages in thread From: Jakub Jelinek @ 2015-11-03 15:38 UTC (permalink / raw) To: Nathan Sidwell; +Cc: GCC Patches On Mon, Nov 02, 2015 at 02:23:19PM -0500, Nathan Sidwell wrote: > Here are the tests for the routine support. The compiler tests check > invalid combinations of gang, worker, vector & seq. The libgomp execution > tests check the expected partioning occurs within loops. As with the > reduction tests, these ones are taken from the execution model loop tests. I find the testsuite coverage insufficient, e.g. you don't have equivalent of first half of declare-simd-2.C or declare-simd-2.c (everything above #pragma omp declare simd inbranch notinbranch), to verify that if acc routine is used without the (fnname) in it, then it can't be followed by var definition and various other tokens. Jakub ^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [2/2] OpenACC routine support 2015-11-03 15:38 ` Jakub Jelinek @ 2015-11-03 15:56 ` Nathan Sidwell 2015-11-03 16:05 ` Jakub Jelinek 0 siblings, 1 reply; 23+ messages in thread From: Nathan Sidwell @ 2015-11-03 15:56 UTC (permalink / raw) To: Jakub Jelinek; +Cc: GCC Patches On 11/03/15 10:38, Jakub Jelinek wrote: > On Mon, Nov 02, 2015 at 02:23:19PM -0500, Nathan Sidwell wrote: >> Here are the tests for the routine support. The compiler tests check >> invalid combinations of gang, worker, vector & seq. The libgomp execution >> tests check the expected partioning occurs within loops. As with the >> reduction tests, these ones are taken from the execution model loop tests. > > I find the testsuite coverage insufficient, e.g. you don't have equivalent > of first half of declare-simd-2.C or declare-simd-2.c > (everything above #pragma omp declare simd inbranch notinbranch), > to verify that if acc routine is used without the (fnname) in it, then > it can't be followed by var definition and various other tokens. d'oh! forgot to port those tests. Easy fix. ok with that added? nathan ^ permalink raw reply [flat|nested] 23+ messages in thread
* Re: [2/2] OpenACC routine support 2015-11-03 15:56 ` Nathan Sidwell @ 2015-11-03 16:05 ` Jakub Jelinek 0 siblings, 0 replies; 23+ messages in thread From: Jakub Jelinek @ 2015-11-03 16:05 UTC (permalink / raw) To: Nathan Sidwell; +Cc: GCC Patches On Tue, Nov 03, 2015 at 10:56:37AM -0500, Nathan Sidwell wrote: > On 11/03/15 10:38, Jakub Jelinek wrote: > >On Mon, Nov 02, 2015 at 02:23:19PM -0500, Nathan Sidwell wrote: > >>Here are the tests for the routine support. The compiler tests check > >>invalid combinations of gang, worker, vector & seq. The libgomp execution > >>tests check the expected partioning occurs within loops. As with the > >>reduction tests, these ones are taken from the execution model loop tests. > > > >I find the testsuite coverage insufficient, e.g. you don't have equivalent > >of first half of declare-simd-2.C or declare-simd-2.c > >(everything above #pragma omp declare simd inbranch notinbranch), > >to verify that if acc routine is used without the (fnname) in it, then > >it can't be followed by var definition and various other tokens. > > d'oh! forgot to port those tests. Easy fix. > > ok with that added? Yes. Jakub ^ permalink raw reply [flat|nested] 23+ messages in thread
end of thread, other threads:[~2015-12-03 14:06 UTC | newest] Thread overview: 23+ messages (download: mbox.gz / follow: Atom feed) -- links below jump to the message on this page -- 2015-11-02 18:56 [0/2] OpenACC routine support Nathan Sidwell 2015-11-02 19:21 ` [1/2] " Nathan Sidwell 2015-11-03 15:35 ` Jakub Jelinek 2015-11-03 15:55 ` Nathan Sidwell 2015-11-10 0:32 ` Nathan Sidwell 2015-11-10 0:48 ` Nathan Sidwell 2015-11-10 5:31 ` Cesar Philippidis 2015-11-10 5:28 ` Cesar Philippidis 2015-11-10 8:16 ` Jakub Jelinek 2015-11-10 14:37 ` Cesar Philippidis 2015-11-18 19:02 ` Cesar Philippidis 2015-11-19 11:23 ` Jakub Jelinek 2015-12-01 14:40 ` Thomas Schwinge 2015-12-01 14:49 ` Cesar Philippidis 2015-12-02 23:37 ` Cesar Philippidis 2015-12-03 8:36 ` Thomas Schwinge 2015-12-03 14:06 ` Cesar Philippidis 2015-11-02 19:23 ` [2/2] " Nathan Sidwell 2015-11-02 19:41 ` Jakub Jelinek 2015-11-02 20:01 ` Nathan Sidwell 2015-11-03 15:38 ` Jakub Jelinek 2015-11-03 15:56 ` Nathan Sidwell 2015-11-03 16:05 ` Jakub Jelinek
This is a public inbox, see mirroring instructions for how to clone and mirror all data and code used for this inbox; as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).