From patchwork Mon Sep 15 23:35:02 2014 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Cesar Philippidis X-Patchwork-Id: 389843 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 7D4E114008F for ; Tue, 16 Sep 2014 09:35:24 +1000 (EST) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender :message-id:date:from:mime-version:to:subject:content-type; q= dns; s=default; b=GAs44VJdB+eodRohmWnF3Q16ZNl/+Si0A4ne1ZzKmCA8GD RcTRqy50P8sEgd2xY7vsOBRMQ39/3ONphncOXB9NpMs0WHvuSO+RLqMbR2Tu+deO xuCsibAOSMwaalvxLNwKUjQpO2o5wpoMITHDcveRQF8VUB3sDm9hUu+nTJa6Y= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender :message-id:date:from:mime-version:to:subject:content-type; s= default; bh=PfrtLie7A/YU4wOQkO1xx9C4wtY=; b=SxuH1bu1VN1JNSsus0z4 zqgimO2Be/IJRKkw9sT8AXJbuk+EKq1gU4xUogYQr6qIsst+yjaJlTZrcZXs7mk3 SAN8XxECnUKDyGs3ktCUc85SanXCFm2Exw+hHTdjAyCOQzgTTIFzl9ZwkCm+Stev 1EZoojUN9StaNbWZUpzoXdg= Received: (qmail 18602 invoked by alias); 15 Sep 2014 23:35:11 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 18580 invoked by uid 89); 15 Sep 2014 23:35:10 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.6 required=5.0 tests=AWL, BAYES_00 autolearn=ham version=3.3.2 X-Spam-User: qpsmtpd, 2 recipients X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Mon, 15 Sep 2014 23:35:07 +0000 Received: from svr-orw-fem-03.mgc.mentorg.com ([147.34.97.39]) by relay1.mentorg.com with esmtp id 1XTfnU-0000Ts-06 from Cesar_Philippidis@mentor.com ; Mon, 15 Sep 2014 16:35:04 -0700 Received: from [127.0.0.1] (147.34.91.1) by svr-orw-fem-03.mgc.mentorg.com (147.34.97.39) with Microsoft SMTP Server id 14.3.181.6; Mon, 15 Sep 2014 16:35:03 -0700 Message-ID: <541777A6.7090908@codesourcery.com> Date: Mon, 15 Sep 2014 16:35:02 -0700 From: Cesar Philippidis User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:31.0) Gecko/20100101 Thunderbird/31.1.1 MIME-Version: 1.0 To: "gcc-patches@gcc.gnu.org" , Subject: [gomp4] OpenACC routine directive This patch adds initial support for the OpenACC routine directive. It's not complete just yet because it doesn't implement any of the optional clauses, except for the optional function/subroutine name. As such, it doesn't go beyond marking functions with the "omp declare target" attribute. According to the OpenACC technical committee, the routine clause will be revised in the next OpenACC 2.x release. In particular, function definitions must have an 'acc routine' associated with it. My understanding is 'acc routine' should also be visible at the call site, but if it's not the compiler can treat it as a regular function call. Furthermore, I've been told that it's not sufficient to place the routine directive in an interface block by itself. E.g. interface recursive function fact (x) !$acc routine integer, intent(in) :: x integer :: fact end function fact end interface integer, parameter :: n = 10 integer :: a(n), i !$acc parallel !$acc loop do i = 1, n a(i) = fact (i) end do !$acc end parallel do i = 1, n write (*, "(I10)") a(i) end do end recursive function fact (x) result (res) integer, intent(in) :: x integer :: res if (x < 1) then res = 1 else res = x * fact (x - 1) end if end function fact This will result in a runtime failure because gcc will not generate an accelerated version of fact. The justification for this is that fortran lacks a file scope, so 'acc routine' wouldn't be visible to fact. Is this patch OK for gomp-4_0-branch? Thanks, Cesar 2014-09-15 Cesar Philippidis gcc/c-family/ * c-pragma.c (oacc_pragmas): Add the "routine" pragma. * c-pragma.h (pragma_kind): Add PRAGMA_OACC_ROUTINE. gcc/c/ * c-parser.c (struct c_parser): New oacc_routines member. (c_parser_external_declaration): Update call to c_parser_declaration_or_fndef. (c_parser_declaration_or_fndef): Add oacc_routine_clauses and oacc_routine_named parameters. Use them in calls to c_finish_oacc_routine and c_parser_declaration_or_fndef. (c_parser_compound_statement_nostart): Update call to c_parser_declaration_or_fndef. (c_parser_label): Likewise. (c_parser_for_statement): Likewise. (c_parser_objc_methodprotolist): Likewise. (c_parser_pragma): Handle PRAGMA_OACC_ROUTINE. (OACC_ROUTINE_CLAUSE_MASK): New macro. (c_parser_oacc_routine): New function. (c_finish_oacc_routine): New function. (c_parser_omp_for_loop): Update calls to c_parser_declaration_or_fndef. (c_parser_omp_declare_simd): Likewise. (c_parse_file): Initialize oacc_routines. gcc/fortran/ * gfortran.h (ST_OACC_ROUTINE): New statement enum. * match.h (gfc_match_oacc_routine): New prototype. * openmp.c (gfc_match_oacc_routine): New function. * parse.c (decode_oacc_directive): Handle the routine directive. (next_statement): Handle ST_OACC_ROUTINE. (gfc_ascii_statement): Likewise. gcc/testsuite/ * c-c++-common/goacc/routine-1.c: New test. * c-c++-common/goacc/routine-2.c: New test. * gfortran.dg/goacc/routine-1.f90: New test. * gfortran.dg/goacc/routine-2.f90: New test. diff --git a/gcc/c-family/c-pragma.c b/gcc/c-family/c-pragma.c index 2d9071a..2ff99f5 100644 --- a/gcc/c-family/c-pragma.c +++ b/gcc/c-family/c-pragma.c @@ -1177,6 +1177,7 @@ static const struct omp_pragma_def oacc_pragmas[] = { { "kernels", PRAGMA_OACC_KERNELS }, { "loop", PRAGMA_OACC_LOOP }, { "parallel", PRAGMA_OACC_PARALLEL }, + { "routine", PRAGMA_OACC_ROUTINE }, { "update", PRAGMA_OACC_UPDATE }, { "wait", PRAGMA_OACC_WAIT }, }; diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h index 4722d51..c097857 100644 --- a/gcc/c-family/c-pragma.h +++ b/gcc/c-family/c-pragma.h @@ -32,6 +32,7 @@ typedef enum pragma_kind { PRAGMA_OACC_KERNELS, PRAGMA_OACC_LOOP, PRAGMA_OACC_PARALLEL, + PRAGMA_OACC_ROUTINE, PRAGMA_OACC_UPDATE, PRAGMA_OACC_WAIT, PRAGMA_OMP_ATOMIC, diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index 09df223..0632c69 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -215,6 +215,10 @@ typedef struct GTY(()) c_parser { /* Buffer to hold all the tokens from parsing the vector attribute for the SIMD-enabled functions (formerly known as elemental functions). */ vec *cilk_simd_fn_tokens; + + /* OpenACC specific parser information. */ + + vec *oacc_routines; } c_parser; @@ -1150,7 +1154,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); + bool, bool, tree *, vec, + tree, bool); 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, @@ -1238,6 +1243,8 @@ static bool c_parser_pragma (c_parser *, enum pragma_context); static bool c_parser_omp_target (c_parser *, enum pragma_context); static void c_parser_omp_end_declare_target (c_parser *); static void c_parser_omp_declare (c_parser *, enum pragma_context); +static void c_parser_oacc_routine (c_parser *parser, enum pragma_context + context); /* These Objective-C parser functions are only ever called when compiling Objective-C. */ @@ -1417,12 +1424,13 @@ c_parser_external_declaration (c_parser *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, false); break; } } static void c_finish_omp_declare_simd (c_parser *, tree, tree, vec); +static void c_finish_oacc_routine (c_parser *, tree, tree, 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 @@ -1500,7 +1508,8 @@ c_parser_declaration_or_fndef (c_parser *parser, bool fndef_ok, bool static_assert_ok, bool empty_ok, bool nested, bool start_attr_ok, tree *objc_foreach_object_declaration, - vec omp_declare_simd_clauses) + vec omp_declare_simd_clauses, + tree oacc_routine_clauses, bool oacc_routine_named) { struct c_declspecs *specs; tree prefix_attrs; @@ -1678,6 +1687,9 @@ c_parser_declaration_or_fndef (c_parser *parser, bool fndef_ok, || !vec_safe_is_empty (parser->cilk_simd_fn_tokens)) c_finish_omp_declare_simd (parser, NULL_TREE, NULL_TREE, omp_declare_simd_clauses); + else + c_finish_oacc_routine (parser, NULL_TREE, + oacc_routine_clauses, oacc_routine_named); c_parser_skip_to_end_of_block_or_statement (parser); return; } @@ -1774,6 +1786,9 @@ c_parser_declaration_or_fndef (c_parser *parser, bool fndef_ok, || !vec_safe_is_empty (parser->cilk_simd_fn_tokens)) c_finish_omp_declare_simd (parser, d, NULL_TREE, omp_declare_simd_clauses); + else + c_finish_oacc_routine (parser, d, oacc_routine_clauses, + oacc_routine_named); } else { @@ -1787,6 +1802,10 @@ c_parser_declaration_or_fndef (c_parser *parser, bool fndef_ok, || !vec_safe_is_empty (parser->cilk_simd_fn_tokens)) c_finish_omp_declare_simd (parser, d, NULL_TREE, omp_declare_simd_clauses); + else + c_finish_oacc_routine (parser, d, oacc_routine_clauses, + oacc_routine_named); + start_init (d, asm_name, global_bindings_p ()); init_loc = c_parser_peek_token (parser)->location; init = c_parser_initializer (parser); @@ -1832,6 +1851,9 @@ c_parser_declaration_or_fndef (c_parser *parser, bool fndef_ok, temp_store_parm_decls (d, parms); c_finish_omp_declare_simd (parser, d, parms, omp_declare_simd_clauses); + c_finish_oacc_routine (parser, d, oacc_routine_clauses, + oacc_routine_named); + if (parms) temp_pop_parm_decls (); } @@ -1938,13 +1960,17 @@ c_parser_declaration_or_fndef (c_parser *parser, bool fndef_ok, function definitions either. */ 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); + c_parser_declaration_or_fndef (parser, false, false, false, true, + false, NULL, vNULL, NULL_TREE, false); 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); + else + c_finish_oacc_routine (parser, current_function_decl, + oacc_routine_clauses, oacc_routine_named); + DECL_STRUCT_FUNCTION (current_function_decl)->function_start_locus = c_parser_peek_token (parser)->location; fnbody = c_parser_compound_statement (parser); @@ -4585,7 +4611,7 @@ c_parser_compound_statement_nostart (c_parser *parser) 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, false); if (last_stmt) pedwarn_c90 (loc, OPT_Wdeclaration_after_statement, "ISO C90 forbids mixed declarations and code"); @@ -4610,7 +4636,8 @@ c_parser_compound_statement_nostart (c_parser *parser) 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, + false); /* Following the old parser, __extension__ does not disable this diagnostic. */ restore_extension_diagnostics (ext); @@ -4747,7 +4774,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, false); } } } @@ -5451,7 +5478,8 @@ c_parser_for_statement (c_parser *parser, bool ivdep) 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, + false); parser->objc_could_be_foreach_context = false; if (c_parser_next_token_is_keyword (parser, RID_IN)) @@ -5480,7 +5508,8 @@ c_parser_for_statement (c_parser *parser, bool ivdep) 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, false); parser->objc_could_be_foreach_context = false; restore_extension_diagnostics (ext); @@ -8568,8 +8597,9 @@ c_parser_objc_methodprotolist (c_parser *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, + false); break; } } @@ -9515,6 +9545,10 @@ c_parser_pragma (c_parser *parser, enum pragma_context context) switch (id) { + case PRAGMA_OACC_ROUTINE: + c_parser_oacc_routine (parser, context); + return false; + case PRAGMA_OACC_UPDATE: if (context != pragma_compound) { @@ -12073,6 +12107,161 @@ c_parser_oacc_parallel (location_t loc, c_parser *parser, char *p_name) return stmt; } +/* OpenACC 2.0: FIXME + # pragma acc routine oacc-routine-clause[optseq] new-line + function-definition + + # pragma acc routine ( name ) oacc-routine-clause[optseq] new-line + + LOC is the location of the #pragma token. +*/ + +#define OACC_ROUTINE_CLAUSE_MASK \ + (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PROC_BIND) + +static void +c_parser_oacc_routine (c_parser *parser, enum pragma_context context) +{ + tree name = NULL_TREE; + location_t here = c_parser_peek_token (parser)->location; + + c_parser_consume_pragma (parser); + + /* Scan for optional '( name )'. */ + if (c_parser_peek_token (parser)->type == CPP_OPEN_PAREN) + { + c_parser_consume_token (parser); + + if (c_parser_next_token_is_not (parser, CPP_NAME) + || c_parser_peek_token (parser)->id_kind != C_ID_ID) + c_parser_error (parser, "expected identifier"); + + // name should be an IDENTIFIER_NODE + name = c_parser_peek_token (parser)->value; + + if (name == NULL_TREE) + { + undeclared_variable (c_parser_peek_token (parser)->location, + c_parser_peek_token (parser)->value); + name = error_mark_node; + } + + c_parser_consume_token (parser); + + if (name == error_mark_node) + return; + + c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, 0); + } + + /* Build a chain of clauses. */ + parser->in_pragma = true; + tree clauses = NULL_TREE; + clauses = c_parser_omp_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK, + "#pragma acc routine"); + + /* Check of the presence if gang, worker, vector and seq clauses, and + throw an error if more than one of those clauses is specified. */ + int parallelism = 0; + tree c; + + for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) + switch (OMP_CLAUSE_CODE (c)) + { + case OMP_CLAUSE_GANG: + case OMP_CLAUSE_WORKER: + case OMP_CLAUSE_VECTOR: + /* FIXME: enable seq. + case OMP_CLAUSE_SEQ: */ + ++parallelism; + break; + default: + break; + } + + if (parallelism > 1) + { + error_at (here, "invalid combination of gang, worker, vector or seq for" + "%<#pragma acc routine%>"); + } + + if (name) + { + TREE_CHAIN (name) = clauses; + vec_safe_push (parser->oacc_routines, name); + } + else + { + if (context != pragma_external) + { + c_parser_error (parser, "%<#pragma acc routine%> must be " + "followed by function declaration or definition"); + return; + } + + if (c_parser_next_token_is (parser, CPP_KEYWORD) + && c_parser_peek_token (parser)->keyword == RID_EXTENSION) + { + int ext = disable_extension_diagnostics (); + do + c_parser_consume_token (parser); + 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, vNULL, clauses, true); + restore_extension_diagnostics (ext); + } + else + c_parser_declaration_or_fndef (parser, true, true, true, false, + true, NULL, vNULL, clauses, true); + } +} + +static void +c_finish_oacc_routine (c_parser *parser, tree fndecl, tree clauses, + bool named) +{ + if (fndecl == NULL_TREE || TREE_CODE (fndecl) != FUNCTION_DECL) + { + if (!named) + return; + + error ("%<#pragma acc routine%> not immediately followed by " + "a function declaration or definition"); + gcc_unreachable(); + return; + } + + if (!named) + { + bool found = false; + int i; + tree t; + + for (i = 0; vec_safe_iterate (parser->oacc_routines, i, &t); i++) + { + if (!strcmp (IDENTIFIER_POINTER (DECL_NAME (fndecl)), + IDENTIFIER_POINTER (t))) + { + found = true; + clauses = TREE_CHAIN (t); + break; + } + } + + if (!found) + return; + } + + if (clauses != NULL_TREE) + clauses = tree_cons (NULL_TREE, clauses, NULL_TREE); + clauses = build_tree_list (get_identifier ("omp declare target"), + clauses); + TREE_CHAIN (clauses) = DECL_ATTRIBUTES (fndecl); + DECL_ATTRIBUTES (fndecl) = clauses; +} + /* OpenACC 2.0: # pragma acc update oacc-update-clause[optseq] new-line */ @@ -12665,7 +12854,7 @@ c_parser_omp_for_loop (location_t loc, c_parser *parser, enum tree_code code, if (i > 0) vec_safe_push (for_block, c_begin_compound_stmt (true)); c_parser_declaration_or_fndef (parser, true, true, true, true, true, - NULL, vNULL); + NULL, vNULL, NULL_TREE, false); decl = check_for_loop_decls (for_loc, flag_isoc99); if (decl == NULL) goto error_init; @@ -13854,12 +14043,12 @@ c_parser_omp_declare_simd (c_parser *parser, enum pragma_context context) 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, false); restore_extension_diagnostics (ext); } else c_parser_declaration_or_fndef (parser, true, true, true, false, true, - NULL, clauses); + NULL, clauses, NULL_TREE, false); break; case pragma_struct: case pragma_param: @@ -13879,7 +14068,8 @@ c_parser_omp_declare_simd (c_parser *parser, enum pragma_context context) 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, + false); restore_extension_diagnostics (ext); break; } @@ -13888,7 +14078,7 @@ c_parser_omp_declare_simd (c_parser *parser, enum pragma_context context) 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, false); break; } c_parser_error (parser, "%<#pragma omp declare simd%> must be followed by " @@ -15230,6 +15420,8 @@ c_parse_file (void) if (tparser.tokens == &tparser.tokens_buf[0]) the_parser->tokens = &the_parser->tokens_buf[0]; + the_parser->oacc_routines = NULL; + /* Initialize EH, if we've been told to do so. */ if (flag_exceptions) using_eh_for_cleanups (); diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h index 755e62e..a379946 100644 --- a/gcc/fortran/gfortran.h +++ b/gcc/fortran/gfortran.h @@ -221,7 +221,7 @@ typedef enum ST_OACC_END_DATA, ST_OACC_HOST_DATA, ST_OACC_END_HOST_DATA, ST_OACC_LOOP, ST_OACC_END_LOOP, ST_OACC_DECLARE, ST_OACC_UPDATE, ST_OACC_WAIT, ST_OACC_CACHE, ST_OACC_KERNELS_LOOP, ST_OACC_END_KERNELS_LOOP, - ST_OACC_ENTER_DATA, ST_OACC_EXIT_DATA, + ST_OACC_ENTER_DATA, ST_OACC_EXIT_DATA, ST_OACC_ROUTINE, ST_OMP_ATOMIC, ST_OMP_BARRIER, ST_OMP_CRITICAL, ST_OMP_END_ATOMIC, ST_OMP_END_CRITICAL, ST_OMP_END_DO, ST_OMP_END_MASTER, ST_OMP_END_ORDERED, ST_OMP_END_PARALLEL, ST_OMP_END_PARALLEL_DO, ST_OMP_END_PARALLEL_SECTIONS, diff --git a/gcc/fortran/match.h b/gcc/fortran/match.h index 3624638..11e199d 100644 --- a/gcc/fortran/match.h +++ b/gcc/fortran/match.h @@ -136,6 +136,7 @@ match gfc_match_oacc_parallel (void); match gfc_match_oacc_parallel_loop (void); match gfc_match_oacc_enter_data (void); match gfc_match_oacc_exit_data (void); +match gfc_match_oacc_routine (void); /* OpenMP directive matchers. */ match gfc_match_omp_eos (void); diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c index ecb20d2..bb3b1b7 100644 --- a/gcc/fortran/openmp.c +++ b/gcc/fortran/openmp.c @@ -1367,6 +1367,79 @@ gfc_match_oacc_cache (void) } +match +gfc_match_oacc_routine (void) +{ + locus old_loc; + gfc_symbol *sym; + match m; + + old_loc = gfc_current_locus; + + m = gfc_match (" ("); + + if (gfc_current_ns->proc_name + && gfc_current_ns->proc_name->attr.if_source == IFSRC_IFBODY + && m == MATCH_YES) + { + gfc_error ("Only the !$ACC ROUTINE form without " + "list is allowed in interface block at %C"); + goto cleanup; + } + + if (m == MATCH_NO + && gfc_current_ns->proc_name + && gfc_match_omp_eos () == MATCH_YES) + { + if (!gfc_add_omp_declare_target (&gfc_current_ns->proc_name->attr, + gfc_current_ns->proc_name->name, + &old_loc)) + goto cleanup; + return MATCH_YES; + } + + if (m != MATCH_YES) + return m; + + /* Scan for a function name. */ + m = gfc_match_symbol (&sym, 0); + + if (m != MATCH_YES) + { + gfc_error ("Syntax error in !$ACC ROUTINE ( NAME ) at %C"); + gfc_current_locus = old_loc; + return MATCH_ERROR; + } + + if (!sym->attr.external && !sym->attr.function && !sym->attr.subroutine) + { + gfc_error ("Syntax error in !$ACC ROUTINE ( NAME ) at %C, invalid" + " function name '%s'", sym->name); + gfc_current_locus = old_loc; + return MATCH_ERROR; + } + + if (gfc_match_char (')') != MATCH_YES) + { + gfc_error ("Syntax error in !$ACC ROUTINE ( NAME ) at %C, expecting" + " ')' after NAME"); + gfc_current_locus = old_loc; + return MATCH_ERROR; + } + + if (gfc_match_omp_eos () != MATCH_YES) + { + gfc_error ("Unexpected junk after !$ACC ROUTINE at %C"); + goto cleanup; + } + return MATCH_YES; + +cleanup: + gfc_current_locus = old_loc; + return MATCH_ERROR; +} + + #define OMP_PARALLEL_CLAUSES \ (OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_SHARED \ | OMP_CLAUSE_COPYIN | OMP_CLAUSE_REDUCTION | OMP_CLAUSE_IF \ diff --git a/gcc/fortran/parse.c b/gcc/fortran/parse.c index b2241d3..411dc27 100644 --- a/gcc/fortran/parse.c +++ b/gcc/fortran/parse.c @@ -644,6 +644,9 @@ decode_oacc_directive (void) case 'l': match ("loop", gfc_match_oacc_loop, ST_OACC_LOOP); break; + case 'r': + match ("routine", gfc_match_oacc_routine, ST_OACC_ROUTINE); + break; case 'u': match ("update", gfc_match_oacc_update, ST_OACC_UPDATE); break; @@ -1356,7 +1359,7 @@ next_statement (void) case ST_EQUIVALENCE: case ST_NAMELIST: case ST_STATEMENT_FUNCTION: \ case ST_TYPE: case ST_INTERFACE: case ST_OMP_THREADPRIVATE: \ case ST_PROCEDURE: case ST_OMP_DECLARE_SIMD: case ST_OMP_DECLARE_REDUCTION: \ - case ST_OMP_DECLARE_TARGET + case ST_OMP_DECLARE_TARGET: case ST_OACC_ROUTINE /* Block end statements. Errors associated with interchanging these are detected in gfc_match_end(). */ @@ -1903,6 +1906,9 @@ gfc_ascii_statement (gfc_statement st) case ST_OACC_EXIT_DATA: p = "!$ACC EXIT DATA"; break; + case ST_OACC_ROUTINE: + p = "!$ACC ROUTINE"; + break; case ST_OMP_ATOMIC: p = "!$OMP ATOMIC"; break; diff --git a/gcc/testsuite/c-c++-common/goacc/routine-1.c b/gcc/testsuite/c-c++-common/goacc/routine-1.c new file mode 100644 index 0000000..1f89fdb --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/routine-1.c @@ -0,0 +1,35 @@ +void *malloc (__SIZE_TYPE__); +void free (void *); + +#pragma acc routine +int +fact (int n) +{ + if (n == 0 || n == 1) + return 1; + + return n * fact (n - 1); +} + +int +main(int argc, char **argv) +{ + int *a, i, n = 10; + + a = (int *)malloc (sizeof (int) * n); + +#pragma acc parallel copy (a[0:n]) vector_length (5) + { +#pragma acc loop + for (i = 0; i < n; i++) + a[i] = fact (i); + } + + for (i = 0; i < n; i++) + if (fact (i) != a[i]) + return -1; + + free (a); + + return 0; +} diff --git a/gcc/testsuite/c-c++-common/goacc/routine-2.c b/gcc/testsuite/c-c++-common/goacc/routine-2.c new file mode 100644 index 0000000..fe2e7f7 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/routine-2.c @@ -0,0 +1,36 @@ +void *malloc (__SIZE_TYPE__); +void free (void *); + +#pragma acc routine (fact) + +int +fact (int n) +{ + if (n == 0 || n == 1) + return 1; + + return n * fact (n - 1); +} + +int +main(int argc, char **argv) +{ + int *a, i, n = 10; + + a = (int *)malloc (sizeof (int) * n); + +#pragma acc parallel copy (a[0:n]) vector_length (5) + { +#pragma acc loop + for (i = 0; i < n; i++) + a[i] = fact (i); + } + + for (i = 0; i < n; i++) + if (fact (i) != a[i]) + return -1; + + free (a); + + return 0; +} diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-1.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-1.f90 new file mode 100644 index 0000000..67c5f11 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/routine-1.f90 @@ -0,0 +1,37 @@ +! { dg-do compile } + + integer, parameter :: n = 10 + integer :: a(n), i + integer, external :: fact + i = 1 + !$acc routine (fact) ! { dg-error "Unexpected \\\!\\\$ACC ROUTINE" } + !$acc routine () ! { dg-error "Syntax error in \\\!\\\$ACC ROUTINE \\\( NAME \\\)" } + !$acc parallel + !$acc loop + do i = 1, n + a(i) = fact (i) + call incr (a(i)) + end do + !$acc end parallel + do i = 1, n + write (*, "(I10)") a(i) + end do +end +recursive function fact (x) result (res) + integer, intent(in) :: x + integer :: res + res = 1 + !$acc routine ! { dg-error "Unexpected \\\!\\\$ACC ROUTINE" } + if (x < 1) then + res = 1 + else + res = x * fact (x - 1) + end if +end function fact +subroutine incr (x) + integer, intent(inout) :: x + integer i + i = 0 + !$acc routine ! { dg-error "Unexpected \\\!\\\$ACC ROUTINE" } + x = x + 1 +end subroutine incr diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-2.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-2.f90 new file mode 100644 index 0000000..3be3351 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/routine-2.f90 @@ -0,0 +1,17 @@ +! { dg-do compile } + + module m1 + contains + recursive function mfact (x) result (res) + integer, intent(in) :: x + integer :: res + integer i + i = 0 + !$acc routine ! { dg-error "Unexpected \\\!\\\$ACC ROUTINE" } + if (x < 1) then + res = 1 + else + res = x * mfact (x - 1) + end if + end function mfact + end module m1