From patchwork Tue Nov 12 09:04:22 2019 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Frederik Harwath X-Patchwork-Id: 1193374 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-513051-incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b="Yw96T4Zx"; dkim-atps=neutral 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 47C1zW3nwcz9s4Y for ; Tue, 12 Nov 2019 20:05:10 +1100 (AEDT) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:date:message-id:in-reply-to:references :mime-version:content-type; q=dns; s=default; b=w+NLSaiIkTWzqp3P PC3n9efb+esSH6CNzqMXOcuV/iy0FXel1/bwGI6p8bc7mnb9YltfxLiIgMtq5J7R pJ4DgK/OBP62CK7DI9WvKEjrKiYcpo+RXUJOFYAJ2VPhWIsyES5VA2Rm9URSoQKE pv9Jugiem4ukiE8QoUfLkvrPzOQ= 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:from :to:cc:subject:date:message-id:in-reply-to:references :mime-version:content-type; s=default; bh=05e2RTF5HTch1vIS1pBeSB ikzJQ=; b=Yw96T4ZxXf5vKGcmSpEQ8VhELftmzYUrmG4+NYzSSLzBQOwvf2tOGw nwGZWMZYoAcoFy6ZWPLxqPOYxhxi9INebq0BrBCo9ldK+6rzZi9eXLIMkvOrA6ba oFRA0bcLGEW9QyYGYjqONl/sTyWMzQ3/cK4jUg/NmtPziW0XtiqFw= Received: (qmail 127305 invoked by alias); 12 Nov 2019 09:04:59 -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 127266 invoked by uid 89); 12 Nov 2019 09:04:59 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-19.2 required=5.0 tests=AWL, BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, SPF_PASS autolearn=ham version=3.3.1 spammy=collapse X-HELO: esa2.mentor.iphmx.com Received: from esa2.mentor.iphmx.com (HELO esa2.mentor.iphmx.com) (68.232.141.98) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Tue, 12 Nov 2019 09:04:49 +0000 IronPort-SDR: 9/rEuapK3/LFEPSU7QxGy8fGlu75BaLimMYVRIUgecvpkfgOGjxgJzmjdr8Q1mwcaKendnxvDC uOUPXbHDz4iT9NaZ+QHy/s5YrD3P0W/TGtqst5AC1ivD4CK7txNfOKpKYWYD6EyPS9tCFAc3xO WNBBtmmvvIlY4hS1zkzKmZN1QBOY3adOswP5i7xr7RRrrlXKv+ZRzfpUvRVAUHhr4gglU3/vqn VDRGXXA/pl9Mzaai8nxxbR500E6Qlx/BBZ3FTNWfl1rCIdeQW5YC5HMelqdriRxwU/JF+lo+ec yLQ= Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa2.mentor.iphmx.com with ESMTP; 12 Nov 2019 01:04:48 -0800 IronPort-SDR: AxsrZSC/g2kueR33guWJfjCVvArOmKpTXp3h9QFV8Mj6kcUqyzovTLGRNi4TXLhMJMBh4IiSLY 9JIm805lTi4aR/YYHU1KxohGH15O95N8+jm+SsQ7JIavSGRSM7SsTOLqBFWhVsDw0/iL81pxyO ezHyZCC0XhtgDWdfsEopD0Qmze/tS3jnjy88+QN7OHgr8NJ3KwAER7EuswpIFdKsVLI/4R96f4 umwNN5yt3uo3m6CeJ25+53DbE4bUvEs+d7Fa8wCVrATQZ+eDyWKi8lMs0IpMJCWC0bxwJMIHPQ gX4= From: To: , , CC: , Subject: [PATCH][COMMITTED] Add OpenACC 2.6 `serial' construct support Date: Tue, 12 Nov 2019 10:04:22 +0100 Message-ID: <20191112090422.7682-1-frederik@codesourcery.com> In-Reply-To: <87pnhyocgd.fsf@euler.schwinge.homeip.net> References: <87pnhyocgd.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 X-IsSubscribed: yes Hi, the following patch has been reviewed and committed. Frederik --------------------------- 8< ------------------------------------------ The `serial' construct (cf. section 2.5.3 of the OpenACC 2.6 standard) is equivalent to a `parallel' construct with clauses `num_gangs(1) num_workers(1) vector_length(1)' implied. These clauses are therefore not supported with the `serial' construct. All the remaining clauses accepted with `parallel' are also accepted with `serial'. The `serial' construct is implemented like `parallel', except for hardcoding dimensions rather than taking them from the relevant clauses, in `expand_omp_target'. Separate codes are used to denote the `serial' construct throughout the middle end, even though the mapping of `serial' to an equivalent `parallel' construct could have been done in the individual language frontends. In particular, this allows to distinguish between compute constructs in warnings, error messages, dumps etc. 2019-11-12 Maciej W. Rozycki Tobias Burnus Frederik Harwath Thomas Schwinge gcc/ * gimple.h (gf_mask): Add GF_OMP_TARGET_KIND_OACC_SERIAL enumeration constant. (is_gimple_omp_oacc): Handle GF_OMP_TARGET_KIND_OACC_SERIAL. (is_gimple_omp_offloaded): Likewise. * gimplify.c (omp_region_type): Add ORT_ACC_SERIAL enumeration constant. Adjust the value of ORT_NONE accordingly. (is_gimple_stmt): Handle OACC_SERIAL. (oacc_default_clause): Handle ORT_ACC_SERIAL. (gomp_needs_data_present): Likewise. (gimplify_adjust_omp_clauses): Likewise. (gimplify_omp_workshare): Handle OACC_SERIAL. (gimplify_expr): Likewise. * omp-expand.c (expand_omp_target): Handle GF_OMP_TARGET_KIND_OACC_SERIAL. (build_omp_regions_1, omp_make_gimple_edges): Likewise. * omp-low.c (is_oacc_parallel): Rename function to... (is_oacc_parallel_or_serial): ... this. Handle GF_OMP_TARGET_KIND_OACC_SERIAL. (scan_sharing_clauses): Adjust accordingly. (scan_omp_for): Likewise. (lower_oacc_head_mark): Likewise. (convert_from_firstprivate_int): Likewise. (lower_omp_target): Likewise. (check_omp_nesting_restrictions): Handle GF_OMP_TARGET_KIND_OACC_SERIAL. (lower_oacc_reductions): Likewise. (lower_omp_target): Likewise. * tree.def (OACC_SERIAL): New tree code. * tree-pretty-print.c (dump_generic_node): Handle OACC_SERIAL. * doc/generic.texi (OpenACC): Document OACC_SERIAL. gcc/c-family/ * c-pragma.h (pragma_kind): Add PRAGMA_OACC_SERIAL enumeration constant. * c-pragma.c (oacc_pragmas): Add "serial" entry. gcc/c/ * c-parser.c (OACC_SERIAL_CLAUSE_MASK): New macro. (c_parser_oacc_kernels_parallel): Rename function to... (c_parser_oacc_compute): ... this. Handle PRAGMA_OACC_SERIAL. (c_parser_omp_construct): Update accordingly. gcc/cp/ * constexpr.c (potential_constant_expression_1): Handle OACC_SERIAL. * parser.c (OACC_SERIAL_CLAUSE_MASK): New macro. (cp_parser_oacc_kernels_parallel): Rename function to... (cp_parser_oacc_compute): ... this. Handle PRAGMA_OACC_SERIAL. (cp_parser_omp_construct): Update accordingly. (cp_parser_pragma): Handle PRAGMA_OACC_SERIAL. Fix alphabetic order. * pt.c (tsubst_expr): Handle OACC_SERIAL. gcc/fortran/ * gfortran.h (gfc_statement): Add ST_OACC_SERIAL_LOOP, ST_OACC_END_SERIAL_LOOP, ST_OACC_SERIAL and ST_OACC_END_SERIAL enumeration constants. (gfc_exec_op): Add EXEC_OACC_SERIAL_LOOP and EXEC_OACC_SERIAL enumeration constants. * match.h (gfc_match_oacc_serial): New prototype. (gfc_match_oacc_serial_loop): Likewise. * dump-parse-tree.c (show_omp_node, show_code_node): Handle EXEC_OACC_SERIAL_LOOP and EXEC_OACC_SERIAL. * match.c (match_exit_cycle): Handle EXEC_OACC_SERIAL_LOOP. * openmp.c (OACC_SERIAL_CLAUSES): New macro. (gfc_match_oacc_serial_loop): New function. (gfc_match_oacc_serial): Likewise. (oacc_is_loop): Handle EXEC_OACC_SERIAL_LOOP. (resolve_omp_clauses): Handle EXEC_OACC_SERIAL. (oacc_code_to_statement): Handle EXEC_OACC_SERIAL and EXEC_OACC_SERIAL_LOOP. (gfc_resolve_oacc_directive): Likewise. * parse.c (decode_oacc_directive) <'s'>: Add case for "serial" and "serial loop". (next_statement): Handle ST_OACC_SERIAL_LOOP and ST_OACC_SERIAL. (gfc_ascii_statement): Likewise. Handle ST_OACC_END_SERIAL_LOOP and ST_OACC_END_SERIAL. (parse_oacc_structured_block): Handle ST_OACC_SERIAL. (parse_oacc_loop): Handle ST_OACC_SERIAL_LOOP and ST_OACC_END_SERIAL_LOOP. (parse_executable): Handle ST_OACC_SERIAL_LOOP and ST_OACC_SERIAL. (is_oacc): Handle EXEC_OACC_SERIAL_LOOP and EXEC_OACC_SERIAL. * resolve.c (gfc_resolve_blocks, gfc_resolve_code): Likewise. * st.c (gfc_free_statement): Likewise. * trans-openmp.c (gfc_trans_oacc_construct): Handle EXEC_OACC_SERIAL. (gfc_trans_oacc_combined_directive): Handle EXEC_OACC_SERIAL_LOOP. (gfc_trans_oacc_directive): Handle EXEC_OACC_SERIAL_LOOP and EXEC_OACC_SERIAL. * trans.c (trans_code): Likewise. gcc/testsuite/ * c-c++-common/goacc/parallel-dims.c: New test. * gfortran.dg/goacc/parallel-dims.f90: New test. libgomp/ * testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: New test. * testsuite/libgomp.oacc-fortran/parallel-dims-aux.c: New test. * testsuite/libgomp.oacc-fortran/parallel-dims.f89: New test. * testsuite/libgomp.oacc-fortran/parallel-dims-2.f90: New test. Reviewed-by: Thomas Schwinge git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@278082 138bc75d-0d04-0410-961f-82ee72b054a4 --- gcc/ChangeLog | 37 ++++++ gcc/c-family/ChangeLog | 8 ++ gcc/c-family/c-pragma.c | 1 + gcc/c-family/c-pragma.h | 1 + gcc/c/ChangeLog | 10 ++ gcc/c/c-parser.c | 34 ++++- gcc/cp/ChangeLog | 14 ++ gcc/cp/constexpr.c | 1 + gcc/cp/parser.c | 35 ++++- gcc/cp/pt.c | 1 + gcc/doc/generic.texi | 5 + gcc/fortran/ChangeLog | 43 +++++++ gcc/fortran/dump-parse-tree.c | 6 + gcc/fortran/gfortran.h | 13 +- gcc/fortran/match.c | 3 +- gcc/fortran/match.h | 2 + gcc/fortran/openmp.c | 33 ++++- gcc/fortran/parse.c | 32 ++++- gcc/fortran/resolve.c | 6 + gcc/fortran/st.c | 2 + gcc/fortran/trans-openmp.c | 17 ++- gcc/fortran/trans.c | 2 + gcc/gimple-pretty-print.c | 3 + gcc/gimple.def | 2 +- gcc/gimple.h | 13 +- gcc/gimplify.c | 22 +++- gcc/omp-expand.c | 43 ++++++- gcc/omp-low.c | 38 ++++-- gcc/testsuite/ChangeLog | 9 ++ .../c-c++-common/goacc/parallel-dims-2.c | 16 ++- .../gfortran.dg/goacc/parallel-dims-2.f90 | 22 ++++ gcc/tree-pretty-print.c | 4 + gcc/tree.def | 6 + gcc/tree.h | 3 +- libgomp/ChangeLog | 10 ++ .../libgomp.oacc-c-c++-common/parallel-dims.c | 73 +++++++++++ .../libgomp.oacc-fortran/parallel-dims-aux.c | 45 +++++++ .../libgomp.oacc-fortran/parallel-dims.f90 | 120 ++++++++++++++++++ 38 files changed, 678 insertions(+), 57 deletions(-) create mode 100644 gcc/testsuite/gfortran.dg/goacc/parallel-dims-2.f90 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/parallel-dims-aux.c create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90 diff --git a/gcc/ChangeLog b/gcc/ChangeLog index b912abcb6135..008e0db21de4 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,40 @@ +2019-11-12 Maciej W. Rozycki + Frederik Harwath + Thomas Schwinge + + gcc/ + * gimple.h (gf_mask): Add GF_OMP_TARGET_KIND_OACC_SERIAL + enumeration constant. + (is_gimple_omp_oacc): Handle GF_OMP_TARGET_KIND_OACC_SERIAL. + (is_gimple_omp_offloaded): Likewise. + * gimplify.c (omp_region_type): Add ORT_ACC_SERIAL enumeration + constant. Adjust the value of ORT_NONE accordingly. + (is_gimple_stmt): Handle OACC_SERIAL. + (oacc_default_clause): Handle ORT_ACC_SERIAL. + (gomp_needs_data_present): Likewise. + (gimplify_adjust_omp_clauses): Likewise. + (gimplify_omp_workshare): Handle OACC_SERIAL. + (gimplify_expr): Likewise. + * omp-expand.c (expand_omp_target): + Handle GF_OMP_TARGET_KIND_OACC_SERIAL. + (build_omp_regions_1, omp_make_gimple_edges): Likewise. + * omp-low.c (is_oacc_parallel): Rename function to... + (is_oacc_parallel_or_serial): ... this. + Handle GF_OMP_TARGET_KIND_OACC_SERIAL. + (scan_sharing_clauses): Adjust accordingly. + (scan_omp_for): Likewise. + (lower_oacc_head_mark): Likewise. + (convert_from_firstprivate_int): Likewise. + (lower_omp_target): Likewise. + (check_omp_nesting_restrictions): Handle + GF_OMP_TARGET_KIND_OACC_SERIAL. + (lower_oacc_reductions): Likewise. + (lower_omp_target): Likewise. + * tree.def (OACC_SERIAL): New tree code. + * tree-pretty-print.c (dump_generic_node): Handle OACC_SERIAL. + + * doc/generic.texi (OpenACC): Document OACC_SERIAL. + 2019-11-12 Jakub Jelinek PR target/92449 diff --git a/gcc/c-family/ChangeLog b/gcc/c-family/ChangeLog index c0efe9005827..3af2c1111981 100644 --- a/gcc/c-family/ChangeLog +++ b/gcc/c-family/ChangeLog @@ -1,3 +1,11 @@ +2019-11-12 Maciej W. Rozycki + Frederik Harwath + + gcc/c-family/ + * c-pragma.h (pragma_kind): Add PRAGMA_OACC_SERIAL enumeration + constant. + * c-pragma.c (oacc_pragmas): Add "serial" entry. + 2019-11-08 Richard Sandiford * c-common.h (gnu_vector_type_p): New function. diff --git a/gcc/c-family/c-pragma.c b/gcc/c-family/c-pragma.c index 9fee84b22383..158154ec1294 100644 --- a/gcc/c-family/c-pragma.c +++ b/gcc/c-family/c-pragma.c @@ -1291,6 +1291,7 @@ static const struct omp_pragma_def oacc_pragmas[] = { { "loop", PRAGMA_OACC_LOOP }, { "parallel", PRAGMA_OACC_PARALLEL }, { "routine", PRAGMA_OACC_ROUTINE }, + { "serial", PRAGMA_OACC_SERIAL }, { "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 e0aa774555a3..bfe681bb430a 100644 --- a/gcc/c-family/c-pragma.h +++ b/gcc/c-family/c-pragma.h @@ -38,6 +38,7 @@ enum pragma_kind { PRAGMA_OACC_LOOP, PRAGMA_OACC_PARALLEL, PRAGMA_OACC_ROUTINE, + PRAGMA_OACC_SERIAL, PRAGMA_OACC_UPDATE, PRAGMA_OACC_WAIT, diff --git a/gcc/c/ChangeLog b/gcc/c/ChangeLog index e3a5120ae98b..a7d86a20c746 100644 --- a/gcc/c/ChangeLog +++ b/gcc/c/ChangeLog @@ -1,3 +1,13 @@ +2019-11-12 Maciej W. Rozycki + Frederik Harwath + + gcc/c/ + * c-parser.c (OACC_SERIAL_CLAUSE_MASK): New macro. + (c_parser_oacc_kernels_parallel): Rename function to... + (c_parser_oacc_compute): ... this. Handle PRAGMA_OACC_SERIAL. + (c_parser_omp_construct): Update accordingly. + + 2019-11-11 Jakub Jelinek * c-parser.c (c_parser_translation_unit): Diagnose declare target diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index 12deb3e76298..8d7ecf400a72 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -16280,6 +16280,11 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name, # pragma acc parallel oacc-parallel-clause[optseq] new-line structured-block + OpenACC 2.6: + + # pragma acc serial oacc-serial-clause[optseq] new-line + structured-block + LOC is the location of the #pragma token. */ @@ -16316,10 +16321,24 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name, | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) +#define OACC_SERIAL_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRIVATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FIRSTPRIVATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) + static tree -c_parser_oacc_kernels_parallel (location_t loc, c_parser *parser, - enum pragma_kind p_kind, char *p_name, - bool *if_p) +c_parser_oacc_compute (location_t loc, c_parser *parser, + enum pragma_kind p_kind, char *p_name, bool *if_p) { omp_clause_mask mask; enum tree_code code; @@ -16335,6 +16354,11 @@ c_parser_oacc_kernels_parallel (location_t loc, c_parser *parser, mask = OACC_PARALLEL_CLAUSE_MASK; code = OACC_PARALLEL; break; + case PRAGMA_OACC_SERIAL: + strcat (p_name, " serial"); + mask = OACC_SERIAL_CLAUSE_MASK; + code = OACC_SERIAL; + break; default: gcc_unreachable (); } @@ -20798,9 +20822,9 @@ c_parser_omp_construct (c_parser *parser, bool *if_p) break; case PRAGMA_OACC_KERNELS: case PRAGMA_OACC_PARALLEL: + case PRAGMA_OACC_SERIAL: strcpy (p_name, "#pragma acc"); - stmt = c_parser_oacc_kernels_parallel (loc, parser, p_kind, p_name, - if_p); + stmt = c_parser_oacc_compute (loc, parser, p_kind, p_name, if_p); break; case PRAGMA_OACC_LOOP: strcpy (p_name, "#pragma acc"); diff --git a/gcc/cp/ChangeLog b/gcc/cp/ChangeLog index 23339b6fc762..972ef791fa7b 100644 --- a/gcc/cp/ChangeLog +++ b/gcc/cp/ChangeLog @@ -1,3 +1,17 @@ +2019-11-12 Maciej W. Rozycki + Frederik Harwath + + gcc/cp/ + * constexpr.c (potential_constant_expression_1): Handle + OACC_SERIAL. + * parser.c (OACC_SERIAL_CLAUSE_MASK): New macro. + (cp_parser_oacc_kernels_parallel): Rename function to... + (cp_parser_oacc_compute): ... this. Handle PRAGMA_OACC_SERIAL. + (cp_parser_omp_construct): Update accordingly. + (cp_parser_pragma): Handle PRAGMA_OACC_SERIAL. Fix alphabetic + order. + * pt.c (tsubst_expr): Handle OACC_SERIAL. + 2019-11-11 Jason Merrill Implement P1946R0, Allow defaulting comparisons by value. diff --git a/gcc/cp/constexpr.c b/gcc/cp/constexpr.c index 20fddc57825a..8c79b0484fce 100644 --- a/gcc/cp/constexpr.c +++ b/gcc/cp/constexpr.c @@ -6986,6 +6986,7 @@ potential_constant_expression_1 (tree t, bool want_rval, bool strict, bool now, case OMP_DEPOBJ: case OACC_PARALLEL: case OACC_KERNELS: + case OACC_SERIAL: case OACC_DATA: case OACC_HOST_DATA: case OACC_LOOP: diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index 1c95d7e9a5af..f2fa7e83952f 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -40175,6 +40175,10 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name, # pragma acc parallel oacc-parallel-clause[optseq] new-line structured-block + + OpenACC 2.6: + + # pragma acc serial oacc-serial-clause[optseq] new-line */ #define OACC_KERNELS_CLAUSE_MASK \ @@ -40210,9 +40214,24 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name, | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) +#define OACC_SERIAL_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRIVATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FIRSTPRIVATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) + static tree -cp_parser_oacc_kernels_parallel (cp_parser *parser, cp_token *pragma_tok, - char *p_name, bool *if_p) +cp_parser_oacc_compute (cp_parser *parser, cp_token *pragma_tok, + char *p_name, bool *if_p) { omp_clause_mask mask; enum tree_code code; @@ -40228,6 +40247,11 @@ cp_parser_oacc_kernels_parallel (cp_parser *parser, cp_token *pragma_tok, mask = OACC_PARALLEL_CLAUSE_MASK; code = OACC_PARALLEL; break; + case PRAGMA_OACC_SERIAL: + strcat (p_name, " serial"); + mask = OACC_SERIAL_CLAUSE_MASK; + code = OACC_SERIAL; + break; default: gcc_unreachable (); } @@ -42047,9 +42071,9 @@ cp_parser_omp_construct (cp_parser *parser, cp_token *pragma_tok, bool *if_p) break; case PRAGMA_OACC_KERNELS: case PRAGMA_OACC_PARALLEL: + case PRAGMA_OACC_SERIAL: strcpy (p_name, "#pragma acc"); - stmt = cp_parser_oacc_kernels_parallel (parser, pragma_tok, p_name, - if_p); + stmt = cp_parser_oacc_compute (parser, pragma_tok, p_name, if_p); break; case PRAGMA_OACC_LOOP: strcpy (p_name, "#pragma acc"); @@ -42716,8 +42740,9 @@ cp_parser_pragma (cp_parser *parser, enum pragma_context context, bool *if_p) case PRAGMA_OACC_DATA: case PRAGMA_OACC_HOST_DATA: case PRAGMA_OACC_KERNELS: - case PRAGMA_OACC_PARALLEL: case PRAGMA_OACC_LOOP: + case PRAGMA_OACC_PARALLEL: + case PRAGMA_OACC_SERIAL: case PRAGMA_OMP_ATOMIC: case PRAGMA_OMP_CRITICAL: case PRAGMA_OMP_DISTRIBUTE: diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c index 8bacb3952ff2..5a0efaa86c8b 100644 --- a/gcc/cp/pt.c +++ b/gcc/cp/pt.c @@ -17991,6 +17991,7 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl, case OACC_KERNELS: case OACC_PARALLEL: + case OACC_SERIAL: tmp = tsubst_omp_clauses (OMP_CLAUSES (t), C_ORT_ACC, args, complain, in_decl); stmt = begin_omp_parallel (); diff --git a/gcc/doc/generic.texi b/gcc/doc/generic.texi index 94e339c15ee8..badaaec38979 100644 --- a/gcc/doc/generic.texi +++ b/gcc/doc/generic.texi @@ -2388,6 +2388,7 @@ compilation. @tindex OACC_KERNELS @tindex OACC_LOOP @tindex OACC_PARALLEL +@tindex OACC_SERIAL @tindex OACC_UPDATE All the statements starting with @code{OACC_} represent directives and @@ -2432,6 +2433,10 @@ See the description of the @code{OMP_FOR} code. Represents @code{#pragma acc parallel [clause1 @dots{} clauseN]}. +@item OACC_SERIAL + +Represents @code{#pragma acc serial [clause1 @dots{} clauseN]}. + @item OACC_UPDATE Represents @code{#pragma acc update [clause1 @dots{} clauseN]}. diff --git a/gcc/fortran/ChangeLog b/gcc/fortran/ChangeLog index 0f14ad1277cf..d92a6ad24804 100644 --- a/gcc/fortran/ChangeLog +++ b/gcc/fortran/ChangeLog @@ -1,3 +1,46 @@ +2019-11-12 Maciej W. Rozycki + Frederik Harwath + + gcc/fortran/ + * gfortran.h (gfc_statement): Add ST_OACC_SERIAL_LOOP, + ST_OACC_END_SERIAL_LOOP, ST_OACC_SERIAL and ST_OACC_END_SERIAL + enumeration constants. + (gfc_exec_op): Add EXEC_OACC_SERIAL_LOOP and EXEC_OACC_SERIAL + enumeration constants. + * match.h (gfc_match_oacc_serial): New prototype. + (gfc_match_oacc_serial_loop): Likewise. + * dump-parse-tree.c (show_omp_node, show_code_node): Handle + EXEC_OACC_SERIAL_LOOP and EXEC_OACC_SERIAL. + * match.c (match_exit_cycle): Handle EXEC_OACC_SERIAL_LOOP. + * openmp.c (OACC_SERIAL_CLAUSES): New macro. + (gfc_match_oacc_serial_loop): New function. + (gfc_match_oacc_serial): Likewise. + (oacc_is_loop): Handle EXEC_OACC_SERIAL_LOOP. + (resolve_omp_clauses): Handle EXEC_OACC_SERIAL. + (oacc_code_to_statement): Handle EXEC_OACC_SERIAL and + EXEC_OACC_SERIAL_LOOP. + (gfc_resolve_oacc_directive): Likewise. + * parse.c (decode_oacc_directive) <'s'>: Add case for "serial" + and "serial loop". + (next_statement): Handle ST_OACC_SERIAL_LOOP and ST_OACC_SERIAL. + (gfc_ascii_statement): Likewise. Handle ST_OACC_END_SERIAL_LOOP + and ST_OACC_END_SERIAL. + (parse_oacc_structured_block): Handle ST_OACC_SERIAL. + (parse_oacc_loop): Handle ST_OACC_SERIAL_LOOP and + ST_OACC_END_SERIAL_LOOP. + (parse_executable): Handle ST_OACC_SERIAL_LOOP and + ST_OACC_SERIAL. + (is_oacc): Handle EXEC_OACC_SERIAL_LOOP and EXEC_OACC_SERIAL. + * resolve.c (gfc_resolve_blocks, gfc_resolve_code): Likewise. + * st.c (gfc_free_statement): Likewise. + * trans-openmp.c (gfc_trans_oacc_construct): Handle + EXEC_OACC_SERIAL. + (gfc_trans_oacc_combined_directive): Handle + EXEC_OACC_SERIAL_LOOP. + (gfc_trans_oacc_directive): Handle EXEC_OACC_SERIAL_LOOP and + EXEC_OACC_SERIAL. + * trans.c (trans_code): Likewise. + 2019-11-11 Janne Blomqvist PR fortran/91828 diff --git a/gcc/fortran/dump-parse-tree.c b/gcc/fortran/dump-parse-tree.c index 9d7aad19e2f5..253fe15b201d 100644 --- a/gcc/fortran/dump-parse-tree.c +++ b/gcc/fortran/dump-parse-tree.c @@ -1654,6 +1654,8 @@ show_omp_node (int level, gfc_code *c) case EXEC_OACC_PARALLEL: name = "PARALLEL"; is_oacc = true; break; case EXEC_OACC_KERNELS_LOOP: name = "KERNELS LOOP"; is_oacc = true; break; case EXEC_OACC_KERNELS: name = "KERNELS"; is_oacc = true; break; + case EXEC_OACC_SERIAL_LOOP: name = "SERIAL LOOP"; is_oacc = true; break; + case EXEC_OACC_SERIAL: name = "SERIAL"; is_oacc = true; break; case EXEC_OACC_DATA: name = "DATA"; is_oacc = true; break; case EXEC_OACC_HOST_DATA: name = "HOST_DATA"; is_oacc = true; break; case EXEC_OACC_LOOP: name = "LOOP"; is_oacc = true; break; @@ -1729,6 +1731,8 @@ show_omp_node (int level, gfc_code *c) case EXEC_OACC_PARALLEL: case EXEC_OACC_KERNELS_LOOP: case EXEC_OACC_KERNELS: + case EXEC_OACC_SERIAL_LOOP: + case EXEC_OACC_SERIAL: case EXEC_OACC_DATA: case EXEC_OACC_HOST_DATA: case EXEC_OACC_LOOP: @@ -2918,6 +2922,8 @@ show_code_node (int level, gfc_code *c) case EXEC_OACC_PARALLEL: case EXEC_OACC_KERNELS_LOOP: case EXEC_OACC_KERNELS: + case EXEC_OACC_SERIAL_LOOP: + case EXEC_OACC_SERIAL: case EXEC_OACC_DATA: case EXEC_OACC_HOST_DATA: case EXEC_OACC_LOOP: diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h index 920acdafc6b7..e962db59bc59 100644 --- a/gcc/fortran/gfortran.h +++ b/gcc/fortran/gfortran.h @@ -223,7 +223,8 @@ enum gfc_statement 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_ROUTINE, + ST_OACC_SERIAL_LOOP, ST_OACC_END_SERIAL_LOOP, ST_OACC_SERIAL, + ST_OACC_END_SERIAL, ST_OACC_ENTER_DATA, ST_OACC_EXIT_DATA, ST_OACC_ROUTINE, ST_OACC_ATOMIC, ST_OACC_END_ATOMIC, 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, @@ -2572,11 +2573,11 @@ enum gfc_exec_op EXEC_BACKSPACE, EXEC_ENDFILE, EXEC_INQUIRE, EXEC_REWIND, EXEC_FLUSH, EXEC_FORM_TEAM, EXEC_CHANGE_TEAM, EXEC_END_TEAM, EXEC_SYNC_TEAM, EXEC_LOCK, EXEC_UNLOCK, EXEC_EVENT_POST, EXEC_EVENT_WAIT, EXEC_FAIL_IMAGE, - EXEC_OACC_KERNELS_LOOP, EXEC_OACC_PARALLEL_LOOP, EXEC_OACC_ROUTINE, - EXEC_OACC_PARALLEL, EXEC_OACC_KERNELS, EXEC_OACC_DATA, EXEC_OACC_HOST_DATA, - EXEC_OACC_LOOP, EXEC_OACC_UPDATE, EXEC_OACC_WAIT, EXEC_OACC_CACHE, - EXEC_OACC_ENTER_DATA, EXEC_OACC_EXIT_DATA, EXEC_OACC_ATOMIC, - EXEC_OACC_DECLARE, + EXEC_OACC_KERNELS_LOOP, EXEC_OACC_PARALLEL_LOOP, EXEC_OACC_SERIAL_LOOP, + EXEC_OACC_ROUTINE, EXEC_OACC_PARALLEL, EXEC_OACC_KERNELS, EXEC_OACC_SERIAL, + EXEC_OACC_DATA, EXEC_OACC_HOST_DATA, EXEC_OACC_LOOP, EXEC_OACC_UPDATE, + EXEC_OACC_WAIT, EXEC_OACC_CACHE, EXEC_OACC_ENTER_DATA, EXEC_OACC_EXIT_DATA, + EXEC_OACC_ATOMIC, EXEC_OACC_DECLARE, EXEC_OMP_CRITICAL, EXEC_OMP_DO, EXEC_OMP_FLUSH, EXEC_OMP_MASTER, EXEC_OMP_ORDERED, EXEC_OMP_PARALLEL, EXEC_OMP_PARALLEL_DO, EXEC_OMP_PARALLEL_SECTIONS, EXEC_OMP_PARALLEL_WORKSHARE, diff --git a/gcc/fortran/match.c b/gcc/fortran/match.c index 4a31080a2856..b5945049de55 100644 --- a/gcc/fortran/match.c +++ b/gcc/fortran/match.c @@ -2860,7 +2860,8 @@ match_exit_cycle (gfc_statement st, gfc_exec_op op) && o != NULL && o->state == COMP_OMP_STRUCTURED_BLOCK && (o->head->op == EXEC_OACC_LOOP - || o->head->op == EXEC_OACC_PARALLEL_LOOP)) + || o->head->op == EXEC_OACC_PARALLEL_LOOP + || o->head->op == EXEC_OACC_SERIAL_LOOP)) { int collapse = 1; gcc_assert (o->head->next != NULL diff --git a/gcc/fortran/match.h b/gcc/fortran/match.h index 611d79646458..7f3d356cbe49 100644 --- a/gcc/fortran/match.h +++ b/gcc/fortran/match.h @@ -146,6 +146,8 @@ match gfc_match_oacc_kernels (void); match gfc_match_oacc_kernels_loop (void); match gfc_match_oacc_parallel (void); match gfc_match_oacc_parallel_loop (void); +match gfc_match_oacc_serial (void); +match gfc_match_oacc_serial_loop (void); match gfc_match_oacc_enter_data (void); match gfc_match_oacc_exit_data (void); match gfc_match_oacc_routine (void); diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c index ca3427885457..dc0521b40f0b 100644 --- a/gcc/fortran/openmp.c +++ b/gcc/fortran/openmp.c @@ -1964,6 +1964,12 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, | OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \ | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_DEFAULT \ | OMP_CLAUSE_WAIT) +#define OACC_SERIAL_CLAUSES \ + (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_REDUCTION \ + | OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \ + | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_DEVICEPTR \ + | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_DEFAULT \ + | OMP_CLAUSE_WAIT) #define OACC_DATA_CLAUSES \ (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_COPY \ | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT | OMP_CLAUSE_CREATE \ @@ -1977,6 +1983,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, (OACC_LOOP_CLAUSES | OACC_PARALLEL_CLAUSES) #define OACC_KERNELS_LOOP_CLAUSES \ (OACC_LOOP_CLAUSES | OACC_KERNELS_CLAUSES) +#define OACC_SERIAL_LOOP_CLAUSES \ + (OACC_LOOP_CLAUSES | OACC_SERIAL_CLAUSES) #define OACC_HOST_DATA_CLAUSES omp_mask (OMP_CLAUSE_USE_DEVICE) #define OACC_DECLARE_CLAUSES \ (omp_mask (OMP_CLAUSE_COPY) | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \ @@ -2038,6 +2046,20 @@ gfc_match_oacc_kernels (void) } +match +gfc_match_oacc_serial_loop (void) +{ + return match_acc (EXEC_OACC_SERIAL_LOOP, OACC_SERIAL_LOOP_CLAUSES); +} + + +match +gfc_match_oacc_serial (void) +{ + return match_acc (EXEC_OACC_SERIAL, OACC_SERIAL_CLAUSES); +} + + match gfc_match_oacc_data (void) { @@ -3783,6 +3805,7 @@ oacc_is_loop (gfc_code *code) { return code->op == EXEC_OACC_PARALLEL_LOOP || code->op == EXEC_OACC_KERNELS_LOOP + || code->op == EXEC_OACC_SERIAL_LOOP || code->op == EXEC_OACC_LOOP; } @@ -4626,7 +4649,9 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, n->sym->name, name, &n->where); } if (code - && (oacc_is_loop (code) || code->op == EXEC_OACC_PARALLEL)) + && (oacc_is_loop (code) + || code->op == EXEC_OACC_PARALLEL + || code->op == EXEC_OACC_SERIAL)) check_array_not_assumed (n->sym, n->where, name); else if (n->sym->as && n->sym->as->type == AS_ASSUMED_SIZE) gfc_error ("Assumed size array %qs in %s clause at %L", @@ -5818,6 +5843,8 @@ oacc_code_to_statement (gfc_code *code) return ST_OACC_PARALLEL; case EXEC_OACC_KERNELS: return ST_OACC_KERNELS; + case EXEC_OACC_SERIAL: + return ST_OACC_SERIAL; case EXEC_OACC_DATA: return ST_OACC_DATA; case EXEC_OACC_HOST_DATA: @@ -5826,6 +5853,8 @@ oacc_code_to_statement (gfc_code *code) return ST_OACC_PARALLEL_LOOP; case EXEC_OACC_KERNELS_LOOP: return ST_OACC_KERNELS_LOOP; + case EXEC_OACC_SERIAL_LOOP: + return ST_OACC_SERIAL_LOOP; case EXEC_OACC_LOOP: return ST_OACC_LOOP; case EXEC_OACC_ATOMIC: @@ -6163,6 +6192,7 @@ gfc_resolve_oacc_directive (gfc_code *code, gfc_namespace *ns ATTRIBUTE_UNUSED) { case EXEC_OACC_PARALLEL: case EXEC_OACC_KERNELS: + case EXEC_OACC_SERIAL: case EXEC_OACC_DATA: case EXEC_OACC_HOST_DATA: case EXEC_OACC_UPDATE: @@ -6174,6 +6204,7 @@ gfc_resolve_oacc_directive (gfc_code *code, gfc_namespace *ns ATTRIBUTE_UNUSED) break; case EXEC_OACC_PARALLEL_LOOP: case EXEC_OACC_KERNELS_LOOP: + case EXEC_OACC_SERIAL_LOOP: case EXEC_OACC_LOOP: resolve_oacc_loop (code); break; diff --git a/gcc/fortran/parse.c b/gcc/fortran/parse.c index 15f6bf2937c4..e44cc6971983 100644 --- a/gcc/fortran/parse.c +++ b/gcc/fortran/parse.c @@ -683,6 +683,9 @@ decode_oacc_directive (void) matcha ("end parallel loop", gfc_match_omp_eos_error, ST_OACC_END_PARALLEL_LOOP); matcha ("end parallel", gfc_match_omp_eos_error, ST_OACC_END_PARALLEL); + matcha ("end serial loop", gfc_match_omp_eos_error, + ST_OACC_END_SERIAL_LOOP); + matcha ("end serial", gfc_match_omp_eos_error, ST_OACC_END_SERIAL); matcha ("enter data", gfc_match_oacc_enter_data, ST_OACC_ENTER_DATA); matcha ("exit data", gfc_match_oacc_exit_data, ST_OACC_EXIT_DATA); break; @@ -705,6 +708,10 @@ decode_oacc_directive (void) case 'r': match ("routine", gfc_match_oacc_routine, ST_OACC_ROUTINE); break; + case 's': + matcha ("serial loop", gfc_match_oacc_serial_loop, ST_OACC_SERIAL_LOOP); + matcha ("serial", gfc_match_oacc_serial, ST_OACC_SERIAL); + break; case 'u': matcha ("update", gfc_match_oacc_update, ST_OACC_UPDATE); break; @@ -1583,7 +1590,8 @@ next_statement (void) case ST_CRITICAL: \ case ST_OACC_PARALLEL_LOOP: case ST_OACC_PARALLEL: case ST_OACC_KERNELS: \ case ST_OACC_DATA: case ST_OACC_HOST_DATA: case ST_OACC_LOOP: \ - case ST_OACC_KERNELS_LOOP: case ST_OACC_ATOMIC + case ST_OACC_KERNELS_LOOP: case ST_OACC_SERIAL_LOOP: case ST_OACC_SERIAL: \ + case ST_OACC_ATOMIC /* Declaration statements */ @@ -2157,6 +2165,18 @@ gfc_ascii_statement (gfc_statement st) case ST_OACC_END_KERNELS_LOOP: p = "!$ACC END KERNELS LOOP"; break; + case ST_OACC_SERIAL_LOOP: + p = "!$ACC SERIAL LOOP"; + break; + case ST_OACC_END_SERIAL_LOOP: + p = "!$ACC END SERIAL LOOP"; + break; + case ST_OACC_SERIAL: + p = "!$ACC SERIAL"; + break; + case ST_OACC_END_SERIAL: + p = "!$ACC END SERIAL"; + break; case ST_OACC_DATA: p = "!$ACC DATA"; break; @@ -5065,6 +5085,9 @@ parse_oacc_structured_block (gfc_statement acc_st) case ST_OACC_KERNELS: acc_end_st = ST_OACC_END_KERNELS; break; + case ST_OACC_SERIAL: + acc_end_st = ST_OACC_END_SERIAL; + break; case ST_OACC_DATA: acc_end_st = ST_OACC_END_DATA; break; @@ -5096,7 +5119,7 @@ parse_oacc_structured_block (gfc_statement acc_st) pop_state (); } -/* Parse the statements of OpenACC loop/parallel loop/kernels loop. */ +/* Parse the statements of OpenACC 'loop', or combined compute 'loop'. */ static gfc_statement parse_oacc_loop (gfc_statement acc_st) @@ -5149,6 +5172,7 @@ parse_oacc_loop (gfc_statement acc_st) gfc_warning (0, "Redundant !$ACC END LOOP at %C"); if ((acc_st == ST_OACC_PARALLEL_LOOP && st == ST_OACC_END_PARALLEL_LOOP) || (acc_st == ST_OACC_KERNELS_LOOP && st == ST_OACC_END_KERNELS_LOOP) || + (acc_st == ST_OACC_SERIAL_LOOP && st == ST_OACC_END_SERIAL_LOOP) || (acc_st == ST_OACC_LOOP && st == ST_OACC_END_LOOP)) { gcc_assert (new_st.op == EXEC_NOP); @@ -5488,6 +5512,7 @@ parse_executable (gfc_statement st) case ST_OACC_PARALLEL_LOOP: case ST_OACC_KERNELS_LOOP: + case ST_OACC_SERIAL_LOOP: case ST_OACC_LOOP: st = parse_oacc_loop (st); if (st == ST_IMPLIED_ENDDO) @@ -5496,6 +5521,7 @@ parse_executable (gfc_statement st) case ST_OACC_PARALLEL: case ST_OACC_KERNELS: + case ST_OACC_SERIAL: case ST_OACC_DATA: case ST_OACC_HOST_DATA: parse_oacc_structured_block (st); @@ -6544,6 +6570,8 @@ is_oacc (gfc_state_data *sd) case EXEC_OACC_PARALLEL: case EXEC_OACC_KERNELS_LOOP: case EXEC_OACC_KERNELS: + case EXEC_OACC_SERIAL_LOOP: + case EXEC_OACC_SERIAL: case EXEC_OACC_DATA: case EXEC_OACC_HOST_DATA: case EXEC_OACC_LOOP: diff --git a/gcc/fortran/resolve.c b/gcc/fortran/resolve.c index a39b9549d7e3..2371a9e201f4 100644 --- a/gcc/fortran/resolve.c +++ b/gcc/fortran/resolve.c @@ -10576,6 +10576,8 @@ gfc_resolve_blocks (gfc_code *b, gfc_namespace *ns) case EXEC_OACC_PARALLEL: case EXEC_OACC_KERNELS_LOOP: case EXEC_OACC_KERNELS: + case EXEC_OACC_SERIAL_LOOP: + case EXEC_OACC_SERIAL: case EXEC_OACC_DATA: case EXEC_OACC_HOST_DATA: case EXEC_OACC_LOOP: @@ -11539,6 +11541,8 @@ gfc_resolve_code (gfc_code *code, gfc_namespace *ns) case EXEC_OACC_PARALLEL: case EXEC_OACC_KERNELS_LOOP: case EXEC_OACC_KERNELS: + case EXEC_OACC_SERIAL_LOOP: + case EXEC_OACC_SERIAL: case EXEC_OACC_DATA: case EXEC_OACC_HOST_DATA: case EXEC_OACC_LOOP: @@ -11952,6 +11956,8 @@ start: case EXEC_OACC_PARALLEL: case EXEC_OACC_KERNELS_LOOP: case EXEC_OACC_KERNELS: + case EXEC_OACC_SERIAL_LOOP: + case EXEC_OACC_SERIAL: case EXEC_OACC_DATA: case EXEC_OACC_HOST_DATA: case EXEC_OACC_LOOP: diff --git a/gcc/fortran/st.c b/gcc/fortran/st.c index ee18d7aea8ad..12eed71e3a26 100644 --- a/gcc/fortran/st.c +++ b/gcc/fortran/st.c @@ -202,6 +202,8 @@ gfc_free_statement (gfc_code *p) case EXEC_OACC_PARALLEL: case EXEC_OACC_KERNELS_LOOP: case EXEC_OACC_KERNELS: + case EXEC_OACC_SERIAL_LOOP: + case EXEC_OACC_SERIAL: case EXEC_OACC_DATA: case EXEC_OACC_HOST_DATA: case EXEC_OACC_LOOP: diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index dee7cc26a7dc..d9dfcabc65ef 100644 --- a/gcc/fortran/trans-openmp.c +++ b/gcc/fortran/trans-openmp.c @@ -3193,8 +3193,9 @@ gfc_trans_omp_code (gfc_code *code, bool force_empty) return stmt; } -/* Trans OpenACC directives. */ -/* parallel, kernels, data and host_data. */ +/* Translate OpenACC 'parallel', 'kernels', 'serial', 'data', 'host_data' + construct. */ + static tree gfc_trans_oacc_construct (gfc_code *code) { @@ -3210,6 +3211,9 @@ gfc_trans_oacc_construct (gfc_code *code) case EXEC_OACC_KERNELS: construct_code = OACC_KERNELS; break; + case EXEC_OACC_SERIAL: + construct_code = OACC_SERIAL; + break; case EXEC_OACC_DATA: construct_code = OACC_DATA; break; @@ -4017,7 +4021,9 @@ gfc_trans_omp_do (gfc_code *code, gfc_exec_op op, stmtblock_t *pblock, return gfc_finish_block (&block); } -/* parallel loop and kernels loop. */ +/* Translate combined OpenACC 'parallel loop', 'kernels loop', 'serial loop' + construct. */ + static tree gfc_trans_oacc_combined_directive (gfc_code *code) { @@ -4035,6 +4041,9 @@ gfc_trans_oacc_combined_directive (gfc_code *code) case EXEC_OACC_KERNELS_LOOP: construct_code = OACC_KERNELS; break; + case EXEC_OACC_SERIAL_LOOP: + construct_code = OACC_SERIAL; + break; default: gcc_unreachable (); } @@ -5267,9 +5276,11 @@ gfc_trans_oacc_directive (gfc_code *code) { case EXEC_OACC_PARALLEL_LOOP: case EXEC_OACC_KERNELS_LOOP: + case EXEC_OACC_SERIAL_LOOP: return gfc_trans_oacc_combined_directive (code); case EXEC_OACC_PARALLEL: case EXEC_OACC_KERNELS: + case EXEC_OACC_SERIAL: case EXEC_OACC_DATA: case EXEC_OACC_HOST_DATA: return gfc_trans_oacc_construct (code); diff --git a/gcc/fortran/trans.c b/gcc/fortran/trans.c index 2f878f6b1185..d9b278199b75 100644 --- a/gcc/fortran/trans.c +++ b/gcc/fortran/trans.c @@ -2137,6 +2137,8 @@ trans_code (gfc_code * code, tree cond) case EXEC_OACC_KERNELS_LOOP: case EXEC_OACC_PARALLEL: case EXEC_OACC_PARALLEL_LOOP: + case EXEC_OACC_SERIAL: + case EXEC_OACC_SERIAL_LOOP: case EXEC_OACC_ENTER_DATA: case EXEC_OACC_EXIT_DATA: case EXEC_OACC_ATOMIC: diff --git a/gcc/gimple-pretty-print.c b/gcc/gimple-pretty-print.c index 2d5ece068053..f59cc2aa3188 100644 --- a/gcc/gimple-pretty-print.c +++ b/gcc/gimple-pretty-print.c @@ -1676,6 +1676,9 @@ dump_gimple_omp_target (pretty_printer *buffer, gomp_target *gs, case GF_OMP_TARGET_KIND_OACC_PARALLEL: kind = " oacc_parallel"; break; + case GF_OMP_TARGET_KIND_OACC_SERIAL: + kind = " oacc_serial"; + break; case GF_OMP_TARGET_KIND_OACC_DATA: kind = " oacc_data"; break; diff --git a/gcc/gimple.def b/gcc/gimple.def index dd64419e8eb6..38c11f41156d 100644 --- a/gcc/gimple.def +++ b/gcc/gimple.def @@ -359,7 +359,7 @@ DEFGSCODE(GIMPLE_OMP_SECTIONS_SWITCH, "gimple_omp_sections_switch", GSS_BASE) DEFGSCODE(GIMPLE_OMP_SINGLE, "gimple_omp_single", GSS_OMP_SINGLE_LAYOUT) /* GIMPLE_OMP_TARGET represents - #pragma acc {kernels,parallel,data,enter data,exit data,update} + #pragma acc {kernels,parallel,serial,data,enter data,exit data,update} #pragma omp target {,data,update} BODY is the sequence of statements inside the construct (NULL for some variants). diff --git a/gcc/gimple.h b/gcc/gimple.h index cf1f8da5ae24..5a190b1714dc 100644 --- a/gcc/gimple.h +++ b/gcc/gimple.h @@ -177,11 +177,12 @@ enum gf_mask { GF_OMP_TARGET_KIND_EXIT_DATA = 4, GF_OMP_TARGET_KIND_OACC_PARALLEL = 5, GF_OMP_TARGET_KIND_OACC_KERNELS = 6, - GF_OMP_TARGET_KIND_OACC_DATA = 7, - GF_OMP_TARGET_KIND_OACC_UPDATE = 8, - GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA = 9, - GF_OMP_TARGET_KIND_OACC_DECLARE = 10, - GF_OMP_TARGET_KIND_OACC_HOST_DATA = 11, + GF_OMP_TARGET_KIND_OACC_SERIAL = 7, + GF_OMP_TARGET_KIND_OACC_DATA = 8, + GF_OMP_TARGET_KIND_OACC_UPDATE = 9, + GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA = 10, + GF_OMP_TARGET_KIND_OACC_DECLARE = 11, + GF_OMP_TARGET_KIND_OACC_HOST_DATA = 12, GF_OMP_TEAMS_GRID_PHONY = 1 << 0, GF_OMP_TEAMS_HOST = 1 << 1, @@ -6476,6 +6477,7 @@ is_gimple_omp_oacc (const gimple *stmt) { case GF_OMP_TARGET_KIND_OACC_PARALLEL: case GF_OMP_TARGET_KIND_OACC_KERNELS: + case GF_OMP_TARGET_KIND_OACC_SERIAL: case GF_OMP_TARGET_KIND_OACC_DATA: case GF_OMP_TARGET_KIND_OACC_UPDATE: case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: @@ -6505,6 +6507,7 @@ is_gimple_omp_offloaded (const gimple *stmt) case GF_OMP_TARGET_KIND_REGION: case GF_OMP_TARGET_KIND_OACC_PARALLEL: case GF_OMP_TARGET_KIND_OACC_KERNELS: + case GF_OMP_TARGET_KIND_OACC_SERIAL: return true; default: return false; diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 2bc41cf98ae9..87a640545141 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -161,6 +161,7 @@ enum omp_region_type ORT_ACC_DATA = ORT_ACC | ORT_TARGET_DATA, /* Data construct. */ ORT_ACC_PARALLEL = ORT_ACC | ORT_TARGET, /* Parallel construct */ ORT_ACC_KERNELS = ORT_ACC | ORT_TARGET | 2, /* Kernels construct. */ + ORT_ACC_SERIAL = ORT_ACC | ORT_TARGET | 4, /* Serial construct. */ ORT_ACC_HOST_DATA = ORT_ACC | ORT_TARGET_DATA | 2, /* Host data. */ /* Dummy OpenMP region, used to disable expansion of @@ -5551,6 +5552,7 @@ is_gimple_stmt (tree t) case STATEMENT_LIST: case OACC_PARALLEL: case OACC_KERNELS: + case OACC_SERIAL: case OACC_DATA: case OACC_HOST_DATA: case OACC_DECLARE: @@ -7289,7 +7291,8 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags) break; case ORT_ACC_PARALLEL: - rkind = "parallel"; + case ORT_ACC_SERIAL: + rkind = ctx->region_type == ORT_ACC_PARALLEL ? "parallel" : "serial"; if (is_private) flags |= GOVD_FIRSTPRIVATE; @@ -10098,10 +10101,11 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, break; } decl = OMP_CLAUSE_DECL (c); - /* Data clauses associated with acc parallel reductions must be + /* Data clauses associated with reductions must be compatible with present_or_copy. Warn and adjust the clause if that is not the case. */ - if (ctx->region_type == ORT_ACC_PARALLEL) + if (ctx->region_type == ORT_ACC_PARALLEL + || ctx->region_type == ORT_ACC_SERIAL) { tree t = DECL_P (decl) ? decl : TREE_OPERAND (decl, 0); n = NULL; @@ -10277,7 +10281,8 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, decl = OMP_CLAUSE_DECL (c); /* OpenACC reductions need a present_or_copy data clause. Add one if necessary. Emit error when the reduction is private. */ - if (ctx->region_type == ORT_ACC_PARALLEL) + if (ctx->region_type == ORT_ACC_PARALLEL + || ctx->region_type == ORT_ACC_SERIAL) { n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl); if (n->value & (GOVD_PRIVATE | GOVD_FIRSTPRIVATE)) @@ -12529,6 +12534,9 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p) case OACC_PARALLEL: ort = ORT_ACC_PARALLEL; break; + case OACC_SERIAL: + ort = ORT_ACC_SERIAL; + break; case OACC_DATA: ort = ORT_ACC_DATA; break; @@ -12612,6 +12620,10 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p) stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_PARALLEL, OMP_CLAUSES (expr)); break; + case OACC_SERIAL: + stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_SERIAL, + OMP_CLAUSES (expr)); + break; case OMP_SECTIONS: stmt = gimple_build_omp_sections (body, OMP_CLAUSES (expr)); break; @@ -13870,6 +13882,7 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p, case OACC_DATA: case OACC_KERNELS: case OACC_PARALLEL: + case OACC_SERIAL: case OMP_SECTIONS: case OMP_SINGLE: case OMP_TARGET: @@ -14286,6 +14299,7 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p, && code != EH_ELSE_EXPR && code != OACC_PARALLEL && code != OACC_KERNELS + && code != OACC_SERIAL && code != OACC_DATA && code != OACC_HOST_DATA && code != OACC_DECLARE diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c index eadff6e50f86..6f945011cf5a 100644 --- a/gcc/omp-expand.c +++ b/gcc/omp-expand.c @@ -7901,12 +7901,14 @@ expand_omp_target (struct omp_region *region) gimple *stmt; edge e; bool offloaded, data_region; + int target_kind; entry_stmt = as_a (last_stmt (region->entry)); + target_kind = gimple_omp_target_kind (entry_stmt); new_bb = region->entry; offloaded = is_gimple_omp_offloaded (entry_stmt); - switch (gimple_omp_target_kind (entry_stmt)) + switch (target_kind) { case GF_OMP_TARGET_KIND_REGION: case GF_OMP_TARGET_KIND_UPDATE: @@ -7914,6 +7916,7 @@ expand_omp_target (struct omp_region *region) case GF_OMP_TARGET_KIND_EXIT_DATA: case GF_OMP_TARGET_KIND_OACC_PARALLEL: case GF_OMP_TARGET_KIND_OACC_KERNELS: + case GF_OMP_TARGET_KIND_OACC_SERIAL: case GF_OMP_TARGET_KIND_OACC_UPDATE: case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: case GF_OMP_TARGET_KIND_OACC_DECLARE: @@ -7944,16 +7947,28 @@ expand_omp_target (struct omp_region *region) entry_bb = region->entry; exit_bb = region->exit; - if (gimple_omp_target_kind (entry_stmt) == GF_OMP_TARGET_KIND_OACC_KERNELS) + switch (target_kind) { + case GF_OMP_TARGET_KIND_OACC_KERNELS: mark_loops_in_oacc_kernels_region (region->entry, region->exit); - /* Further down, both OpenACC kernels and OpenACC parallel constructs - will be mappted to BUILT_IN_GOACC_PARALLEL, and to distinguish the - two, there is an "oacc kernels" attribute set for OpenACC kernels. */ + /* Further down, all OpenACC compute constructs will be mapped to + BUILT_IN_GOACC_PARALLEL, and to distinguish between them, there + is an "oacc kernels" attribute set for OpenACC kernels. */ DECL_ATTRIBUTES (child_fn) = tree_cons (get_identifier ("oacc kernels"), NULL_TREE, DECL_ATTRIBUTES (child_fn)); + break; + case GF_OMP_TARGET_KIND_OACC_SERIAL: + /* Further down, all OpenACC compute constructs will be mapped to + BUILT_IN_GOACC_PARALLEL, and to distinguish between them, there + is an "oacc serial" attribute set for OpenACC serial. */ + DECL_ATTRIBUTES (child_fn) + = tree_cons (get_identifier ("oacc serial"), + NULL_TREE, DECL_ATTRIBUTES (child_fn)); + break; + default: + break; } if (offloaded) @@ -8156,8 +8171,9 @@ expand_omp_target (struct omp_region *region) start_ix = BUILT_IN_GOMP_TARGET_ENTER_EXIT_DATA; flags_i |= GOMP_TARGET_FLAG_EXIT_DATA; break; - case GF_OMP_TARGET_KIND_OACC_KERNELS: case GF_OMP_TARGET_KIND_OACC_PARALLEL: + case GF_OMP_TARGET_KIND_OACC_KERNELS: + case GF_OMP_TARGET_KIND_OACC_SERIAL: start_ix = BUILT_IN_GOACC_PARALLEL; break; case GF_OMP_TARGET_KIND_OACC_DATA: @@ -8352,7 +8368,18 @@ expand_omp_target (struct omp_region *region) args.quick_push (get_target_arguments (&gsi, entry_stmt)); break; case BUILT_IN_GOACC_PARALLEL: - oacc_set_fn_attrib (child_fn, clauses, &args); + if (lookup_attribute ("oacc serial", DECL_ATTRIBUTES (child_fn)) != NULL) + { + tree dims = NULL_TREE; + unsigned int ix; + + /* For serial constructs we set all dimensions to 1. */ + for (ix = GOMP_DIM_MAX; ix--;) + dims = tree_cons (NULL_TREE, integer_one_node, dims); + oacc_replace_fn_attrib (child_fn, dims); + } + else + oacc_set_fn_attrib (child_fn, clauses, &args); tagging = true; /* FALLTHRU */ case BUILT_IN_GOACC_ENTER_EXIT_DATA: @@ -8913,6 +8940,7 @@ build_omp_regions_1 (basic_block bb, struct omp_region *parent, case GF_OMP_TARGET_KIND_DATA: case GF_OMP_TARGET_KIND_OACC_PARALLEL: case GF_OMP_TARGET_KIND_OACC_KERNELS: + case GF_OMP_TARGET_KIND_OACC_SERIAL: case GF_OMP_TARGET_KIND_OACC_DATA: case GF_OMP_TARGET_KIND_OACC_HOST_DATA: break; @@ -9167,6 +9195,7 @@ omp_make_gimple_edges (basic_block bb, struct omp_region **region, case GF_OMP_TARGET_KIND_DATA: case GF_OMP_TARGET_KIND_OACC_PARALLEL: case GF_OMP_TARGET_KIND_OACC_KERNELS: + case GF_OMP_TARGET_KIND_OACC_SERIAL: case GF_OMP_TARGET_KIND_OACC_DATA: case GF_OMP_TARGET_KIND_OACC_HOST_DATA: break; diff --git a/gcc/omp-low.c b/gcc/omp-low.c index e232d7aa62d3..781e7cbf27a2 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -185,15 +185,18 @@ static tree scan_omp_1_op (tree *, int *, void *); *handled_ops_p = false; \ break; -/* Return true if CTX corresponds to an oacc parallel region. */ +/* Return true if CTX corresponds to an OpenACC 'parallel' or 'serial' + region. */ static bool -is_oacc_parallel (omp_context *ctx) +is_oacc_parallel_or_serial (omp_context *ctx) { enum gimple_code outer_type = gimple_code (ctx->stmt); return ((outer_type == GIMPLE_OMP_TARGET) - && (gimple_omp_target_kind (ctx->stmt) - == GF_OMP_TARGET_KIND_OACC_PARALLEL)); + && ((gimple_omp_target_kind (ctx->stmt) + == GF_OMP_TARGET_KIND_OACC_PARALLEL) + || (gimple_omp_target_kind (ctx->stmt) + == GF_OMP_TARGET_KIND_OACC_SERIAL))); } /* Return true if CTX corresponds to an oacc kernels region. */ @@ -1149,7 +1152,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) goto do_private; case OMP_CLAUSE_REDUCTION: - if (is_oacc_parallel (ctx) || is_oacc_kernels (ctx)) + if (is_oacc_parallel_or_serial (ctx) || is_oacc_kernels (ctx)) ctx->local_reduction_clauses = tree_cons (NULL, c, ctx->local_reduction_clauses); /* FALLTHRU */ @@ -2391,7 +2394,7 @@ scan_omp_for (gomp_for *stmt, omp_context *outer_ctx) { omp_context *tgt = enclosing_target_ctx (outer_ctx); - if (!tgt || is_oacc_parallel (tgt)) + if (!tgt || is_oacc_parallel_or_serial (tgt)) for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) { char const *check = NULL; @@ -2417,7 +2420,7 @@ scan_omp_for (gomp_for *stmt, omp_context *outer_ctx) if (check && OMP_CLAUSE_OPERAND (c, 0)) error_at (gimple_location (stmt), "argument not permitted on %qs clause in" - " OpenACC %", check); + " OpenACC % or %", check); } if (tgt && is_oacc_kernels (tgt)) @@ -2945,6 +2948,7 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx) { case GF_OMP_TARGET_KIND_OACC_PARALLEL: case GF_OMP_TARGET_KIND_OACC_KERNELS: + case GF_OMP_TARGET_KIND_OACC_SERIAL: ok = true; break; @@ -3393,6 +3397,7 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx) stmt_name = "target exit data"; break; case GF_OMP_TARGET_KIND_OACC_PARALLEL: stmt_name = "parallel"; break; case GF_OMP_TARGET_KIND_OACC_KERNELS: stmt_name = "kernels"; break; + case GF_OMP_TARGET_KIND_OACC_SERIAL: stmt_name = "serial"; break; case GF_OMP_TARGET_KIND_OACC_DATA: stmt_name = "data"; break; case GF_OMP_TARGET_KIND_OACC_UPDATE: stmt_name = "update"; break; case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: @@ -3410,6 +3415,8 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx) ctx_stmt_name = "parallel"; break; case GF_OMP_TARGET_KIND_OACC_KERNELS: ctx_stmt_name = "kernels"; break; + case GF_OMP_TARGET_KIND_OACC_SERIAL: + ctx_stmt_name = "serial"; break; case GF_OMP_TARGET_KIND_OACC_DATA: ctx_stmt_name = "data"; break; case GF_OMP_TARGET_KIND_OACC_HOST_DATA: ctx_stmt_name = "host_data"; break; @@ -6711,8 +6718,10 @@ lower_oacc_reductions (location_t loc, tree clauses, tree level, bool inner, break; case GIMPLE_OMP_TARGET: - if (gimple_omp_target_kind (probe->stmt) - != GF_OMP_TARGET_KIND_OACC_PARALLEL) + if ((gimple_omp_target_kind (probe->stmt) + != GF_OMP_TARGET_KIND_OACC_PARALLEL) + && (gimple_omp_target_kind (probe->stmt) + != GF_OMP_TARGET_KIND_OACC_SERIAL)) goto do_lookup; cls = gimple_omp_target_clauses (probe->stmt); @@ -7518,7 +7527,7 @@ lower_oacc_head_mark (location_t loc, tree ddvar, tree clauses, /* In a parallel region, loops are implicitly INDEPENDENT. */ omp_context *tgt = enclosing_target_ctx (ctx); - if (!tgt || is_oacc_parallel (tgt)) + if (!tgt || is_oacc_parallel_or_serial (tgt)) tag |= OLF_INDEPENDENT; if (tag & OLF_TILE) @@ -11357,6 +11366,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) case GF_OMP_TARGET_KIND_EXIT_DATA: case GF_OMP_TARGET_KIND_OACC_PARALLEL: case GF_OMP_TARGET_KIND_OACC_KERNELS: + case GF_OMP_TARGET_KIND_OACC_SERIAL: case GF_OMP_TARGET_KIND_OACC_UPDATE: case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: case GF_OMP_TARGET_KIND_OACC_DECLARE: @@ -11489,7 +11499,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) if (!maybe_lookup_field (var, ctx)) continue; - /* Don't remap oacc parallel reduction variables, because the + /* Don't remap compute constructs' reduction variables, because the intermediate result must be local to each gang. */ if (offloaded && !(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && OMP_CLAUSE_MAP_IN_REDUCTION (c))) @@ -11531,7 +11541,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) break; case OMP_CLAUSE_FIRSTPRIVATE: - if (is_oacc_parallel (ctx)) + if (is_oacc_parallel_or_serial (ctx)) goto oacc_firstprivate; map_cnt++; var = OMP_CLAUSE_DECL (c); @@ -11905,7 +11915,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) break; case OMP_CLAUSE_FIRSTPRIVATE: - if (is_oacc_parallel (ctx)) + if (is_oacc_parallel_or_serial (ctx)) goto oacc_firstprivate_map; ovar = OMP_CLAUSE_DECL (c); if (omp_is_reference (ovar)) @@ -12509,7 +12519,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) gimple_seq fork_seq = NULL; gimple_seq join_seq = NULL; - if (is_oacc_parallel (ctx)) + if (is_oacc_parallel_or_serial (ctx)) { /* If there are reductions on the offloaded region itself, treat them as a dummy GANG loop. */ diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog index 2196bf0d5019..45fb8e520184 100644 --- a/gcc/testsuite/ChangeLog +++ b/gcc/testsuite/ChangeLog @@ -1,3 +1,12 @@ +2019-11-12 Maciej W. Rozycki + Tobias Burnus + Frederik Harwath + Thomas Schwinge + + gcc/testsuite/ + * c-c++-common/goacc/parallel-dims.c: New test. + * gfortran.dg/goacc/parallel-dims.f90: New test. + 2019-11-12 Jakub Jelinek PR tree-optimization/92452 diff --git a/gcc/testsuite/c-c++-common/goacc/parallel-dims-2.c b/gcc/testsuite/c-c++-common/goacc/parallel-dims-2.c index acfbe7ff031a..31c4ee349f2c 100644 --- a/gcc/testsuite/c-c++-common/goacc/parallel-dims-2.c +++ b/gcc/testsuite/c-c++-common/goacc/parallel-dims-2.c @@ -1,5 +1,7 @@ -/* Invalid use of OpenACC parallelism dimensions clauses: num_gangs, - num_workers, vector_length. */ +/* Invalid use of OpenACC parallelism dimensions clauses: 'num_gangs', + 'num_workers', 'vector_length'. */ + +/* See also '../../gfortran.dg/goacc/parallel-dims-2.f90'. */ void f(int i, float f) { @@ -255,4 +257,14 @@ void f(int i, float f) vector_length(&f) /* { dg-error "'vector_length' expression must be integral" } */ \ num_gangs( /* { dg-error "expected (primary-|)expression before end of line" "TODO" { xfail c } } */ ; + + + /* The 'serial' construct doesn't allow these at all. */ + +#pragma acc serial num_gangs (1) /* { dg-error "'num_gangs' is not valid for '#pragma acc serial'" } */ + ; +#pragma acc serial num_workers (1) /* { dg-error "'num_workers' is not valid for '#pragma acc serial'" } */ + ; +#pragma acc serial vector_length (1) /* { dg-error "'vector_length' is not valid for '#pragma acc serial'" } */ + ; } diff --git a/gcc/testsuite/gfortran.dg/goacc/parallel-dims-2.f90 b/gcc/testsuite/gfortran.dg/goacc/parallel-dims-2.f90 new file mode 100644 index 000000000000..91a5c300a94c --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/parallel-dims-2.f90 @@ -0,0 +1,22 @@ +! Invalid use of OpenACC parallelism dimensions clauses: 'num_gangs', +! 'num_workers', 'vector_length'. + +! See also '../../c-c++-common/goacc/parallel-dims-2.c'. + +subroutine f() + !TODO 'kernels', 'parallel' testing per '../../c-c++-common/goacc/parallel-dims-2.c'. + !TODO This should incorporate some of the testing done in 'sie.f95'. + + + ! The 'serial' construct doesn't allow these at all. + +!$acc serial num_gangs (1) ! { dg-error "Failed to match clause at" } +!$acc end serial ! { dg-error "Unexpected !.ACC END SERIAL statement" } + +!$acc serial num_workers (1) ! { dg-error "Failed to match clause at" } +!$acc end serial ! { dg-error "Unexpected !.ACC END SERIAL statement" } + +!$acc serial vector_length (1) ! { dg-error "Failed to match clause at" } +!$acc end serial ! { dg-error "Unexpected !.ACC END SERIAL statement" } + +end subroutine f diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c index 53b3f55a3e6a..1cf7a9121336 100644 --- a/gcc/tree-pretty-print.c +++ b/gcc/tree-pretty-print.c @@ -3223,6 +3223,10 @@ dump_generic_node (pretty_printer *pp, tree node, int spc, dump_flags_t flags, pp_string (pp, "#pragma acc kernels"); goto dump_omp_clauses_body; + case OACC_SERIAL: + pp_string (pp, "#pragma acc serial"); + goto dump_omp_clauses_body; + case OACC_DATA: pp_string (pp, "#pragma acc data"); dump_omp_clauses (pp, OACC_DATA_CLAUSES (node), spc, flags); diff --git a/gcc/tree.def b/gcc/tree.def index fb6e7344fa6b..e8bb4f37f802 100644 --- a/gcc/tree.def +++ b/gcc/tree.def @@ -1095,6 +1095,12 @@ DEFTREECODE (OACC_PARALLEL, "oacc_parallel", tcc_statement, 2) DEFTREECODE (OACC_KERNELS, "oacc_kernels", tcc_statement, 2) +/* OpenACC - #pragma acc serial [clause1 ... clauseN] + Operand 0: OMP_BODY: Code to be executed sequentially. + Operand 1: OMP_CLAUSES: List of clauses. */ + +DEFTREECODE (OACC_SERIAL, "oacc_serial", tcc_statement, 2) + /* OpenACC - #pragma acc data [clause1 ... clauseN] Operand 0: OACC_DATA_BODY: Data construct body. Operand 1: OACC_DATA_CLAUSES: List of clauses. */ diff --git a/gcc/tree.h b/gcc/tree.h index a7d39c3a74df..4bec90d9a729 100644 --- a/gcc/tree.h +++ b/gcc/tree.h @@ -1622,7 +1622,8 @@ class auto_suppress_location_wrappers treatment if OMP_CLAUSE_SIZE is zero. */ #define OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION(NODE) \ TREE_PROTECTED (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)) -/* Nonzero if this map clause is for an ACC parallel reduction variable. */ +/* Nonzero if this map clause is for an OpenACC compute construct's reduction + variable. */ #define OMP_CLAUSE_MAP_IN_REDUCTION(NODE) \ TREE_PRIVATE (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)) diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 2f60d606a88c..734395936f05 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,13 @@ +2019-11-12 Maciej W. Rozycki + Tobias Burnus + Frederik Harwath + Thomas Schwinge + + libgomp/ + * testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: New test. + * testsuite/libgomp.oacc-fortran/parallel-dims-aux.c: New test. + * testsuite/libgomp.oacc-fortran/parallel-dims.f89: New test. + 2019-11-11 Tobias Burnus Kwok Cheung Yeung diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c index 7e699f476b21..a5edfc6ca164 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c @@ -1,6 +1,8 @@ /* OpenACC parallelism dimensions clauses: num_gangs, num_workers, vector_length. */ +/* See also '../libgomp.oacc-fortran/parallel-dims.f90'. */ + #include #include #include @@ -45,6 +47,8 @@ int main () { acc_init (acc_device_default); + /* OpenACC parallel construct. */ + /* Non-positive value. */ /* GR, WS, VS. */ @@ -478,6 +482,8 @@ int main () } + /* OpenACC kernels construct. */ + /* We can't test parallelized OpenACC kernels constructs in this way: use of the acc_gang, acc_worker, acc_vector functions will make the construct unparallelizable. */ @@ -544,5 +550,72 @@ int main () } + /* OpenACC serial construct. */ + + /* GR, WS, VS. */ + { + int gangs_min, gangs_max, workers_min, workers_max, vectors_min, vectors_max; + gangs_min = workers_min = vectors_min = INT_MAX; + gangs_max = workers_max = vectors_max = INT_MIN; +#pragma acc serial /* { dg-warning "using vector_length \$32\$, ignoring 1" "" { target openacc_nvidia_accel_selected } } */ \ + reduction (min: gangs_min, workers_min, vectors_min) reduction (max: gangs_max, workers_max, vectors_max) + { + for (int i = 100; i > -100; i--) + { + gangs_min = gangs_max = acc_gang (); + workers_min = workers_max = acc_worker (); + vectors_min = vectors_max = acc_vector (); + } + } + if (gangs_min != 0 || gangs_max != 1 - 1 + || workers_min != 0 || workers_max != 1 - 1 + || vectors_min != 0 || vectors_max != 1 - 1) + __builtin_abort (); + } + + /* Composition of GP, WP, VP. */ + { + int vectors_actual = 1; /* Implicit 'vector_length (1)' clause. */ + int gangs_min, gangs_max, workers_min, workers_max, vectors_min, vectors_max; + gangs_min = workers_min = vectors_min = INT_MAX; + gangs_max = workers_max = vectors_max = INT_MIN; +#pragma acc serial copy (vectors_actual) /* { dg-warning "using vector_length \$32\$, ignoring 1" "" { target openacc_nvidia_accel_selected } } */ \ + copy (gangs_min, gangs_max, workers_min, workers_max, vectors_min, vectors_max) + { + if (acc_on_device (acc_device_nvidia)) + { + /* The GCC nvptx back end enforces vector_length (32). */ + /* It's unclear if that's actually permissible here; + "OpenACC + 'serial' construct might not actually be serial". */ + vectors_actual = 32; + } +#pragma acc loop gang reduction (min: gangs_min, workers_min, vectors_min) reduction (max: gangs_max, workers_max, vectors_max) + for (int i = 100; i > -100; i--) +#pragma acc loop worker reduction (min: gangs_min, workers_min, vectors_min) reduction (max: gangs_max, workers_max, vectors_max) + for (int j = 100; j > -100; j--) +#pragma acc loop vector reduction (min: gangs_min, workers_min, vectors_min) reduction (max: gangs_max, workers_max, vectors_max) + for (int k = 100 * vectors_actual; k > -100 * vectors_actual; k--) + { + gangs_min = gangs_max = acc_gang (); + workers_min = workers_max = acc_worker (); + vectors_min = vectors_max = acc_vector (); + } + } + if (acc_get_device_type () == acc_device_nvidia) + { + if (vectors_actual != 32) + __builtin_abort (); + } + else + if (vectors_actual != 1) + __builtin_abort (); + if (gangs_min != 0 || gangs_max != 1 - 1 + || workers_min != 0 || workers_max != 1 - 1 + || vectors_min != 0 || vectors_max != vectors_actual - 1) + __builtin_abort (); + } + + return 0; } diff --git a/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims-aux.c b/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims-aux.c new file mode 100644 index 000000000000..b5986f4afef7 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims-aux.c @@ -0,0 +1,45 @@ +/* OpenACC parallelism dimensions clauses: num_gangs, num_workers, + vector_length. */ + +/* Copied from '../libgomp.oacc-c-c++-common/parallel-dims.c'. */ + +/* Used by 'parallel-dims.f90'. */ + +#include +#include +#include + +/* TODO: "(int) acc_device_*" casts because of the C++ acc_on_device wrapper + not behaving as expected for -O0. */ +#pragma acc routine seq +/* static */ unsigned int __attribute__ ((optimize ("O2"))) acc_gang () +{ + if (acc_on_device ((int) acc_device_host)) + return 0; + else if (acc_on_device ((int) acc_device_nvidia)) + return __builtin_goacc_parlevel_id (GOMP_DIM_GANG); + else + __builtin_abort (); +} + +#pragma acc routine seq +/* static */ unsigned int __attribute__ ((optimize ("O2"))) acc_worker () +{ + if (acc_on_device ((int) acc_device_host)) + return 0; + else if (acc_on_device ((int) acc_device_nvidia)) + return __builtin_goacc_parlevel_id (GOMP_DIM_WORKER); + else + __builtin_abort (); +} + +#pragma acc routine seq +/* static */ unsigned int __attribute__ ((optimize ("O2"))) acc_vector () +{ + if (acc_on_device ((int) acc_device_host)) + return 0; + else if (acc_on_device ((int) acc_device_nvidia)) + return __builtin_goacc_parlevel_id (GOMP_DIM_VECTOR); + else + __builtin_abort (); +} diff --git a/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90 b/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90 new file mode 100644 index 000000000000..1bfcd6ce0998 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90 @@ -0,0 +1,120 @@ +! OpenACC parallelism dimensions clauses: num_gangs, num_workers, +! vector_length. + +! { dg-additional-sources parallel-dims-aux.c } +! { dg-do run } +! { dg-prune-output "command-line option '-fintrinsic-modules-path=.*' is valid for Fortran but not for C" } + +! See also '../libgomp.oacc-c-c++-common/parallel-dims.c'. + +module acc_routines + implicit none (type, external) + + interface + integer function acc_gang() bind(C) + !$acc routine seq + end function acc_gang + + integer function acc_worker() bind(C) + !$acc routine seq + end function acc_worker + + integer function acc_vector() bind(C) + !$acc routine seq + end function acc_vector + end interface +end module acc_routines + +program main + use iso_c_binding + use openacc + use acc_routines + implicit none (type, external) + + integer :: gangs_min, gangs_max, workers_min, workers_max, vectors_min, vectors_max + integer :: vectors_actual + integer :: i, j, k + + call acc_init (acc_device_default) + + ! OpenACC parallel construct. + + !TODO + + + ! OpenACC kernels construct. + + !TODO + + + ! OpenACC serial construct. + + ! GR, WS, VS. + + gangs_min = huge(gangs_min) ! INT_MAX + workers_min = huge(workers_min) ! INT_MAX + vectors_min = huge(vectors_min) ! INT_MAX + gangs_max = -huge(gangs_max) - 1 ! INT_MIN + workers_max = -huge(gangs_max) - 1 ! INT_MIN + vectors_max = -huge(gangs_max) - 1 ! INT_MIN + !$acc serial & + !$acc reduction (min: gangs_min, workers_min, vectors_min) reduction (max: gangs_max, workers_max, vectors_max) ! { dg-warning "using vector_length \$32\$, ignoring 1" "" { target openacc_nvidia_accel_selected } } + do i = 100, -99, -1 + gangs_min = acc_gang (); + gangs_max = acc_gang (); + workers_min = acc_worker (); + workers_max = acc_worker (); + vectors_min = acc_vector (); + vectors_max = acc_vector (); + end do + !$acc end serial + if (gangs_min /= 0 .or. gangs_max /= 1 - 1 & + .or. workers_min /= 0 .or. workers_max /= 1 - 1 & + .or. vectors_min /= 0 .or. vectors_max /= 1 - 1) & + stop 1 + + ! Composition of GP, WP, VP. + + vectors_actual = 1 ! Implicit 'vector_length (1)' clause. + gangs_min = huge(gangs_min) ! INT_MAX + workers_min = huge(workers_min) ! INT_MAX + vectors_min = huge(vectors_min) ! INT_MAX + gangs_max = -huge(gangs_max) - 1 ! INT_MIN + workers_max = -huge(gangs_max) - 1 ! INT_MIN + vectors_max = -huge(gangs_max) - 1 ! INT_MIN + !$acc serial copy (vectors_actual) & + !$acc copy (gangs_min, gangs_max, workers_min, workers_max, vectors_min, vectors_max) ! { dg-warning "using vector_length \$32\$, ignoring 1" "" { target openacc_nvidia_accel_selected } } + if (acc_on_device (acc_device_nvidia)) then + ! The GCC nvptx back end enforces vector_length (32). + ! It's unclear if that's actually permissible here; + ! "OpenACC 'serial' + ! construct might not actually be serial". + vectors_actual = 32 + end if + !$acc loop gang reduction (min: gangs_min, workers_min, vectors_min) reduction (max: gangs_max, workers_max, vectors_max) + do i = 100, -99, -1 + !$acc loop worker reduction (min: gangs_min, workers_min, vectors_min) reduction (max: gangs_max, workers_max, vectors_max) + do j = 100, -99, -1 + !$acc loop vector reduction (min: gangs_min, workers_min, vectors_min) reduction (max: gangs_max, workers_max, vectors_max) + do k = 100 * vectors_actual, -99 * vectors_actual, -1 + gangs_min = acc_gang (); + gangs_max = acc_gang (); + workers_min = acc_worker (); + workers_max = acc_worker (); + vectors_min = acc_vector (); + vectors_max = acc_vector (); + end do + end do + end do + !$acc end serial + if (acc_get_device_type () .eq. acc_device_nvidia) then + if (vectors_actual /= 32) stop 2 + else + if (vectors_actual /= 1) stop 3 + end if + if (gangs_min /= 0 .or. gangs_max /= 1 - 1 & + .or. workers_min /= 0 .or. workers_max /= 1 - 1 & + .or. vectors_min /= 0 .or. vectors_max /= vectors_actual - 1) & + stop 4 + +end program main