From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 36255 invoked by alias); 17 Dec 2018 11:10:28 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Received: (qmail 36147 invoked by uid 89); 17 Dec 2018 11:10:27 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-4.2 required=5.0 tests=BAYES_50,GIT_PATCH_3,RCVD_IN_DNSWL_NONE,SPF_PASS,TIME_LIMIT_EXCEEDED autolearn=unavailable version=3.3.2 spammy=consequently, worker, promoted, seq 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, 17 Dec 2018 11:10:16 +0000 Received: from svr-orw-mbx-02.mgc.mentorg.com ([147.34.90.202]) by relay1.mentorg.com with esmtps (TLSv1.2:ECDHE-RSA-AES256-SHA384:256) id 1gYqms-0002bo-KF from ChungLin_Tang@mentor.com ; Mon, 17 Dec 2018 03:10:14 -0800 Received: from [0.0.0.0] (147.34.91.1) by svr-orw-mbx-02.mgc.mentorg.com (147.34.90.202) with Microsoft SMTP Server (TLS) id 15.0.1320.4; Mon, 17 Dec 2018 03:10:09 -0800 Reply-To: Subject: Re: [PATCH, og8] Add OpenACC 2.6 `serial' construct support To: "Maciej W. Rozycki" , Joseph Myers , CC: Thomas Schwinge , Chung-Lin Tang , Jakub Jelinek , Catherine Moore References: From: Chung-Lin Tang Message-ID: <6dc98006-0de4-d831-55e1-6f538c0734fc@mentor.com> Date: Mon, 17 Dec 2018 11:10:00 -0000 User-Agent: Mozilla/5.0 (Macintosh; Intel Mac OS X 10.13; rv:60.0) Gecko/20100101 Thunderbird/60.3.3 MIME-Version: 1.0 In-Reply-To: Content-Type: text/plain; charset="utf-8"; format=flowed Content-Transfer-Encoding: 7bit X-SW-Source: 2018-12/txt/msg01200.txt.bz2 Hi Maciej, I don't think there's anything wrong with the patch itself; if the testcases pass, then it should implement the functionality correctly. My only issue is should "serial" really be promoted to such a visible construct in the middle-end? It's just a special case of parallel, and user debug errors can be dealt with specifically. I don't see much value of it being preserved past the front-ends into gimplify/omp-low, just more testing to be done to guard various conditions that are specific to OpenACC... That's just my 0.02, others can provide more input. Chung-Lin On 2018/12/17 11:09 AM, Maciej W. Rozycki wrote: > The `serial' construct is equivalent to a `parallel' construct with > clauses `num_gangs(1) num_workers(1) vector_length(1)' implied. > Naturally these clauses are therefore not supported with the `serial' > construct. All the remaining clauses accepted with `parallel' are also > accepted with `serial'. > > Consequently implementation is straightforward, by handling `serial' > exactly 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, saving a lot of mechanical changes and avoiding middle-end > code expansion. This is so that any reporting such as with warning or > error messages and in diagnostic dumps use `serial' rather than > `parallel', therefore avoiding user confusion. > > 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. > (omp_add_variable): Handle ORT_ACC_SERIAL. > (oacc_default_clause): Likewise. > (gimplify_scan_omp_clauses): Likewise. > (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. > (build_receiver_ref): Adjust accordingly. > (build_sender_ref): Likewise. > (scan_sharing_clauses): Likewise. > (create_omp_child_function): Likewise. > (scan_omp_for): Likewise. > (scan_omp_target): 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-pretty-print.c (dump_generic_node): Handle OACC_SERIAL. > * tree.def (OACC_SERIAL): New tree code. > > * 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. > (OACC_SERIAL_CLAUSE_DEVICE_TYPE_MASK): Likewise. > (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. > (OACC_SERIAL_CLAUSE_DEVICE_TYPE_MASK): Likewise. > (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. > (OACC_SERIAL_CLAUSE_DEVICE_TYPE_MASK): Likewise. > (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_is_serial): New function. > (oacc_code_to_statement): Handle EXEC_OACC_SERIAL and > EXEC_OACC_SERIAL_LOOP. > (gfc_resolve_oacc_directive): Likewise. > (resolve_oacc_loop_blocks): Also call `oacc_is_serial'. > * 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/serial-dims.c: New test. > > libgomp/ > * testsuite/libgomp.oacc-c-c++-common/serial-dims.c: New test. > --- > Hi, > > I find the: > > if ((ctx->region_type & (ORT_ACC_PARALLEL | ORT_ACC_KERNELS)) != 0 > > statement near the beginning of `oacc_default_clause' highly suspicious > and unfortunately it was added with r230275 with no discussion (cf. > ). AFAICT > syntactically it amounts to: > > if ((ctx->region_type & ORT_ACC_KERNELS) != 0 > > (because ORT_ACC_KERNELS is bitwise a superset of ORT_ACC_PARALLEL) or: > > if ((ctx->region_type & (ORT_ACC | ORT_TARGET | 0x80)) != 0 > > which already covers ORT_ACC_SERIAL, so I have decided not to add it here. > Furthermore `oacc_default_clause' is only ever called when ORT_ACC is set: > > if ((ctx->region_type & ORT_ACC) != 0) > nflags = oacc_default_clause (ctx, decl, flags); > > so that condition actually always evaluates to true. > > Perhaps: > > if ((ctx->region_type == ORT_ACC_PARALLEL > || ctx->region_type == ORT_ACC_KERNELS) > > was meant instead, in which case ORT_ACC_SERIAL would have to be listed > explicitly, but I would be wary of blindly changing code that has been out > there for 3 years now and obviously must have worked, without having a > test case to verify such a change. > > Joseph, you are listed as a co-author of r230275: is that a piece of that > change you would be able to comment on by any chance? > > This has passed regression-testing with the `x86_64-linux-gnu' target and > the `nvptx-none' offload target, across the `gcc', `g++', `gfortran' and > `libgomp' test suites. I will appreciate feedback and if none has been > given in a couple of days' time, then I will commit this change to the og8 > branch. > > A Fortran test case equivalent to C/C++ `serial-dims.c' would be good > having, but Fortran programming has not been my strongest skill and I > didn't want to delay this submission. I'll see if I can make one before > the final commit. > > Maciej > --- > gcc/c-family/c-pragma.c | 1 > gcc/c-family/c-pragma.h | 1 > gcc/c/c-parser.c | 41 +++++ > gcc/cp/constexpr.c | 1 > gcc/cp/parser.c | 42 +++++- > gcc/cp/pt.c | 1 > gcc/doc/generic.texi | 5 > 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 | 52 +++++++ > gcc/fortran/parse.c | 27 +++ > gcc/fortran/resolve.c | 6 > gcc/fortran/st.c | 2 > gcc/fortran/trans-openmp.c | 14 +- > gcc/fortran/trans.c | 2 > gcc/gimple-pretty-print.c | 3 > gcc/gimple.h | 3 > gcc/gimplify.c | 35 +++-- > gcc/omp-expand.c | 37 ++++- > gcc/omp-low.c | 56 ++++---- > gcc/testsuite/c-c++-common/goacc/serial-dims.c | 12 + > gcc/tree-pretty-print.c | 4 > gcc/tree.def | 6 > libgomp/testsuite/libgomp.oacc-c-c++-common/serial-dims.c | 98 ++++++++++++++ > 26 files changed, 415 insertions(+), 58 deletions(-) > > gcc-openacc-serial.diff > Index: gcc-openacc-gcc-8-branch/gcc/c-family/c-pragma.c > =================================================================== > --- gcc-openacc-gcc-8-branch.orig/gcc/c-family/c-pragma.c > +++ gcc-openacc-gcc-8-branch/gcc/c-family/c-pragma.c > @@ -1277,6 +1277,7 @@ static const struct omp_pragma_def oacc_ > { "loop", PRAGMA_OACC_LOOP }, > { "parallel", PRAGMA_OACC_PARALLEL }, > { "routine", PRAGMA_OACC_ROUTINE }, > + { "serial", PRAGMA_OACC_SERIAL }, > { "update", PRAGMA_OACC_UPDATE }, > { "wait", PRAGMA_OACC_WAIT } > }; > Index: gcc-openacc-gcc-8-branch/gcc/c-family/c-pragma.h > =================================================================== > --- gcc-openacc-gcc-8-branch.orig/gcc/c-family/c-pragma.h > +++ gcc-openacc-gcc-8-branch/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, > > Index: gcc-openacc-gcc-8-branch/gcc/c/c-parser.c > =================================================================== > --- gcc-openacc-gcc-8-branch.orig/gcc/c/c-parser.c > +++ gcc-openacc-gcc-8-branch/gcc/c/c-parser.c > @@ -14949,6 +14949,11 @@ c_parser_oacc_loop (location_t loc, c_pa > # 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. > */ > > @@ -15003,6 +15008,27 @@ c_parser_oacc_loop (location_t loc, c_pa > | (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_ATTACH) \ > + | (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_DEVICE_TYPE) \ > + | (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) ) > + > +#define OACC_SERIAL_CLAUSE_DEVICE_TYPE_MASK \ > + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \ > + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) > + > static tree > mark_vars_oacc_gangprivate (tree *tp, > int *walk_subtrees ATTRIBUTE_UNUSED, > @@ -15031,9 +15057,8 @@ mark_vars_oacc_gangprivate (tree *tp, > } > > 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, dmask; > enum tree_code code; > @@ -15051,6 +15076,12 @@ c_parser_oacc_kernels_parallel (location > dmask = OACC_PARALLEL_CLAUSE_DEVICE_TYPE_MASK; > code = OACC_PARALLEL; > break; > + case PRAGMA_OACC_SERIAL: > + strcat (p_name, " serial"); > + mask = OACC_SERIAL_CLAUSE_MASK; > + dmask = OACC_SERIAL_CLAUSE_DEVICE_TYPE_MASK; > + code = OACC_SERIAL; > + break; > default: > gcc_unreachable (); > } > @@ -18347,9 +18378,9 @@ c_parser_omp_construct (c_parser *parser > 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"); > Index: gcc-openacc-gcc-8-branch/gcc/cp/constexpr.c > =================================================================== > --- gcc-openacc-gcc-8-branch.orig/gcc/cp/constexpr.c > +++ gcc-openacc-gcc-8-branch/gcc/cp/constexpr.c > @@ -5690,6 +5690,7 @@ potential_constant_expression_1 (tree t, > case OMP_ATOMIC_CAPTURE_NEW: > case OACC_PARALLEL: > case OACC_KERNELS: > + case OACC_SERIAL: > case OACC_DATA: > case OACC_HOST_DATA: > case OACC_LOOP: > Index: gcc-openacc-gcc-8-branch/gcc/cp/parser.c > =================================================================== > --- gcc-openacc-gcc-8-branch.orig/gcc/cp/parser.c > +++ gcc-openacc-gcc-8-branch/gcc/cp/parser.c > @@ -37255,6 +37255,10 @@ cp_parser_oacc_loop (cp_parser *parser, > > # 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 \ > @@ -37308,6 +37312,27 @@ cp_parser_oacc_loop (cp_parser *parser, > | (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_ATTACH) \ > + | (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_DEVICE_TYPE) \ > + | (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) ) > + > +#define OACC_SERIAL_CLAUSE_DEVICE_TYPE_MASK \ > + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \ > + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) > + > tree > mark_vars_oacc_gangprivate (tree *tp, > int *walk_subtrees ATTRIBUTE_UNUSED, > @@ -37337,8 +37362,8 @@ mark_vars_oacc_gangprivate (tree *tp, > } > > 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, dmask; > enum tree_code code; > @@ -37356,6 +37381,12 @@ cp_parser_oacc_kernels_parallel (cp_pars > dmask = OACC_PARALLEL_CLAUSE_DEVICE_TYPE_MASK; > code = OACC_PARALLEL; > break; > + case PRAGMA_OACC_SERIAL: > + strcat (p_name, " serial"); > + mask = OACC_SERIAL_CLAUSE_MASK; > + dmask = OACC_SERIAL_CLAUSE_DEVICE_TYPE_MASK; > + code = OACC_SERIAL; > + break; > default: > gcc_unreachable (); > } > @@ -38550,9 +38581,9 @@ cp_parser_omp_construct (cp_parser *pars > 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"); > @@ -39187,8 +39218,9 @@ cp_parser_pragma (cp_parser *parser, enu > 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: > Index: gcc-openacc-gcc-8-branch/gcc/cp/pt.c > =================================================================== > --- gcc-openacc-gcc-8-branch.orig/gcc/cp/pt.c > +++ gcc-openacc-gcc-8-branch/gcc/cp/pt.c > @@ -17074,6 +17074,7 @@ tsubst_expr (tree t, tree args, tsubst_f > > 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 (); > Index: gcc-openacc-gcc-8-branch/gcc/doc/generic.texi > =================================================================== > --- gcc-openacc-gcc-8-branch.orig/gcc/doc/generic.texi > +++ gcc-openacc-gcc-8-branch/gcc/doc/generic.texi > @@ -2355,6 +2355,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 > @@ -2399,6 +2400,10 @@ See the description of the @code{OMP_FOR > > 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]}. > Index: gcc-openacc-gcc-8-branch/gcc/fortran/dump-parse-tree.c > =================================================================== > --- gcc-openacc-gcc-8-branch.orig/gcc/fortran/dump-parse-tree.c > +++ gcc-openacc-gcc-8-branch/gcc/fortran/dump-parse-tree.c > @@ -1538,6 +1538,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; > @@ -1613,6 +1615,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: > @@ -2798,6 +2802,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: > Index: gcc-openacc-gcc-8-branch/gcc/fortran/gfortran.h > =================================================================== > --- gcc-openacc-gcc-8-branch.orig/gcc/fortran/gfortran.h > +++ gcc-openacc-gcc-8-branch/gcc/fortran/gfortran.h > @@ -222,7 +222,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, > @@ -2524,11 +2525,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, > Index: gcc-openacc-gcc-8-branch/gcc/fortran/match.c > =================================================================== > --- gcc-openacc-gcc-8-branch.orig/gcc/fortran/match.c > +++ gcc-openacc-gcc-8-branch/gcc/fortran/match.c > @@ -2817,7 +2817,8 @@ match_exit_cycle (gfc_statement st, gfc_ > && 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 > Index: gcc-openacc-gcc-8-branch/gcc/fortran/match.h > =================================================================== > --- gcc-openacc-gcc-8-branch.orig/gcc/fortran/match.h > +++ gcc-openacc-gcc-8-branch/gcc/fortran/match.h > @@ -146,6 +146,8 @@ match gfc_match_oacc_kernels_loop (void) > match gfc_match_oacc_parallel (void); > match gfc_match_oacc_parallel_loop (void); > match gfc_match_oacc_enter_data (void); > +match gfc_match_oacc_serial (void); > +match gfc_match_oacc_serial_loop (void); > match gfc_match_oacc_exit_data (void); > match gfc_match_oacc_routine (void); > > Index: gcc-openacc-gcc-8-branch/gcc/fortran/openmp.c > =================================================================== > --- gcc-openacc-gcc-8-branch.orig/gcc/fortran/openmp.c > +++ gcc-openacc-gcc-8-branch/gcc/fortran/openmp.c > @@ -2084,6 +2084,16 @@ gfc_match_omp_clauses (gfc_omp_clauses * > | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT \ > | OMP_CLAUSE_DEVICEPTR \ > | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_ATTACH) > +#define OACC_SERIAL_CLAUSES \ > + (omp_mask (OMP_CLAUSE_ASYNC) | OMP_CLAUSE_WAIT \ > + | OMP_CLAUSE_DEVICE_TYPE \ > + | OMP_CLAUSE_IF \ > + | 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_ATTACH) > #define OACC_DATA_CLAUSES \ > (omp_mask (OMP_CLAUSE_IF) \ > | OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \ > @@ -2141,6 +2151,9 @@ gfc_match_omp_clauses (gfc_omp_clauses * > | OMP_CLAUSE_NUM_GANGS | OMP_CLAUSE_NUM_WORKERS \ > | OMP_CLAUSE_VECTOR_LENGTH \ > | OMP_CLAUSE_DEVICE_TYPE) > +#define OACC_SERIAL_CLAUSE_DEVICE_TYPE_MASK \ > + (omp_mask (OMP_CLAUSE_ASYNC) | OMP_CLAUSE_WAIT \ > + | OMP_CLAUSE_DEVICE_TYPE) > #define OACC_LOOP_CLAUSE_DEVICE_TYPE_MASK \ > (omp_mask (OMP_CLAUSE_COLLAPSE) \ > | OMP_CLAUSE_GANG | OMP_CLAUSE_WORKER | OMP_CLAUSE_VECTOR \ > @@ -2207,6 +2220,24 @@ gfc_match_oacc_kernels (void) > > > match > +gfc_match_oacc_serial_loop (void) > +{ > + return match_acc (EXEC_OACC_SERIAL_LOOP, > + OACC_SERIAL_CLAUSES | OACC_LOOP_CLAUSES, > + OACC_SERIAL_CLAUSE_DEVICE_TYPE_MASK > + | OACC_LOOP_CLAUSE_DEVICE_TYPE_MASK); > +} > + > + > +match > +gfc_match_oacc_serial (void) > +{ > + return match_acc (EXEC_OACC_SERIAL, OACC_SERIAL_CLAUSES, > + OACC_SERIAL_CLAUSE_DEVICE_TYPE_MASK); > +} > + > + > +match > gfc_match_oacc_data (void) > { > return match_acc (EXEC_OACC_DATA, OACC_DATA_CLAUSES, OMP_MASK2_LAST); > @@ -3995,6 +4026,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; > } > > @@ -4807,7 +4839,9 @@ resolve_omp_clauses (gfc_code *code, gfc > 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", > @@ -5968,6 +6002,12 @@ oacc_is_kernels (gfc_code *code) > return code->op == EXEC_OACC_KERNELS || code->op == EXEC_OACC_KERNELS_LOOP; > } > > +static bool > +oacc_is_serial (gfc_code *code) > +{ > + return code->op == EXEC_OACC_SERIAL || code->op == EXEC_OACC_SERIAL_LOOP; > +} > + > static gfc_statement > omp_code_to_statement (gfc_code *code) > { > @@ -6009,6 +6049,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: > @@ -6017,6 +6059,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: > @@ -6198,7 +6242,9 @@ resolve_oacc_loop_blocks (gfc_code *code > &code->loc); > } > > - if (oacc_is_parallel (c->code) || oacc_is_kernels (c->code)) > + if (oacc_is_parallel (c->code) > + || oacc_is_kernels (c->code) > + || oacc_is_serial (c->code)) > break; > } > > @@ -6415,6 +6461,7 @@ gfc_resolve_oacc_directive (gfc_code *co > { > 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: > @@ -6426,6 +6473,7 @@ gfc_resolve_oacc_directive (gfc_code *co > 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; > Index: gcc-openacc-gcc-8-branch/gcc/fortran/parse.c > =================================================================== > --- gcc-openacc-gcc-8-branch.orig/gcc/fortran/parse.c > +++ gcc-openacc-gcc-8-branch/gcc/fortran/parse.c > @@ -690,6 +690,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; > @@ -1541,7 +1545,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 */ > > @@ -2109,6 +2114,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; > @@ -4927,6 +4944,9 @@ parse_oacc_structured_block (gfc_stateme > 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; > @@ -5011,6 +5031,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); > @@ -5346,6 +5367,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) > @@ -5354,6 +5376,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); > @@ -6346,6 +6369,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: > Index: gcc-openacc-gcc-8-branch/gcc/fortran/resolve.c > =================================================================== > --- gcc-openacc-gcc-8-branch.orig/gcc/fortran/resolve.c > +++ gcc-openacc-gcc-8-branch/gcc/fortran/resolve.c > @@ -10090,6 +10090,8 @@ gfc_resolve_blocks (gfc_code *b, gfc_nam > 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: > @@ -11037,6 +11039,8 @@ gfc_resolve_code (gfc_code *code, gfc_na > 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: > @@ -11445,6 +11449,8 @@ gfc_resolve_code (gfc_code *code, gfc_na > 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: > Index: gcc-openacc-gcc-8-branch/gcc/fortran/st.c > =================================================================== > --- gcc-openacc-gcc-8-branch.orig/gcc/fortran/st.c > +++ gcc-openacc-gcc-8-branch/gcc/fortran/st.c > @@ -201,6 +201,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: > Index: gcc-openacc-gcc-8-branch/gcc/fortran/trans-openmp.c > =================================================================== > --- gcc-openacc-gcc-8-branch.orig/gcc/fortran/trans-openmp.c > +++ gcc-openacc-gcc-8-branch/gcc/fortran/trans-openmp.c > @@ -3305,7 +3305,7 @@ gfc_init_nodesc_arrays (stmtblock_t *inn > } > > /* Trans OpenACC directives. */ > -/* parallel, kernels, data and host_data. */ > +/* parallel, serial, kernels, data and host_data. */ > static tree > gfc_trans_oacc_construct (gfc_code *code) > { > @@ -3325,6 +3325,10 @@ gfc_trans_oacc_construct (gfc_code *code > construct_code = OACC_KERNELS; > scan_nodesc_arrays = true; > break; > + case EXEC_OACC_SERIAL: > + construct_code = OACC_SERIAL; > + scan_nodesc_arrays = true; > + break; > case EXEC_OACC_DATA: > construct_code = OACC_DATA; > break; > @@ -4210,7 +4214,7 @@ gfc_filter_oacc_combined_clauses (gfc_om > construct_code); > } > > -/* Combined OpenACC parallel loop and kernels loop. */ > +/* Combined OpenACC parallel loop, kernels loop and serial loop. */ > static tree > gfc_trans_oacc_combined_directive (gfc_code *code) > { > @@ -4232,6 +4236,10 @@ gfc_trans_oacc_combined_directive (gfc_c > construct_code = OACC_KERNELS; > scan_nodesc_arrays = true; > break; > + case EXEC_OACC_SERIAL_LOOP: > + construct_code = OACC_SERIAL; > + scan_nodesc_arrays = true; > + break; > default: > gcc_unreachable (); > } > @@ -5480,9 +5488,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); > Index: gcc-openacc-gcc-8-branch/gcc/fortran/trans.c > =================================================================== > --- gcc-openacc-gcc-8-branch.orig/gcc/fortran/trans.c > +++ gcc-openacc-gcc-8-branch/gcc/fortran/trans.c > @@ -2109,6 +2109,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: > Index: gcc-openacc-gcc-8-branch/gcc/gimple-pretty-print.c > =================================================================== > --- gcc-openacc-gcc-8-branch.orig/gcc/gimple-pretty-print.c > +++ gcc-openacc-gcc-8-branch/gcc/gimple-pretty-print.c > @@ -1605,6 +1605,9 @@ dump_gimple_omp_target (pretty_printer * > 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; > Index: gcc-openacc-gcc-8-branch/gcc/gimple.h > =================================================================== > --- gcc-openacc-gcc-8-branch.orig/gcc/gimple.h > +++ gcc-openacc-gcc-8-branch/gcc/gimple.h > @@ -183,6 +183,7 @@ enum gf_mask { > 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 = 12, > GF_OMP_TEAMS_GRID_PHONY = 1 << 0, > > /* True on an GIMPLE_OMP_RETURN statement if the return does not require > @@ -6299,6 +6300,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: > @@ -6328,6 +6330,7 @@ is_gimple_omp_offloaded (const gimple *s > 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; > Index: gcc-openacc-gcc-8-branch/gcc/gimplify.c > =================================================================== > --- gcc-openacc-gcc-8-branch.orig/gcc/gimplify.c > +++ gcc-openacc-gcc-8-branch/gcc/gimplify.c > @@ -147,11 +147,12 @@ 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 | 0x80, /* Kernels construct. */ > + ORT_ACC_SERIAL = ORT_ACC | ORT_TARGET | 0x100, /* Serial construct. */ > ORT_ACC_HOST_DATA = ORT_ACC | ORT_TARGET_DATA | 0x80, /* Host data. */ > > /* Dummy OpenMP region, used to disable expansion of > DECL_VALUE_EXPRs in taskloop pre body. */ > - ORT_NONE = 0x100 > + ORT_NONE = 0x200 > }; > > /* Gimplify hashtable helper. */ > @@ -5450,6 +5451,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: > @@ -6947,7 +6949,8 @@ omp_add_variable (struct gimplify_omp_ct > map_private = oacc_privatize_reduction (ctx->outer_context); > > if (ctx->outer_context > - && ctx->outer_context->region_type == ORT_ACC_PARALLEL) > + && (ctx->outer_context->region_type == ORT_ACC_PARALLEL > + || ctx->outer_context->region_type == ORT_ACC_SERIAL)) > update_data_map = true; > > while (outer_ctx) > @@ -6967,7 +6970,8 @@ omp_add_variable (struct gimplify_omp_ct > && (n->value & GOVD_MAP)); > } > else if (update_data_map > - && outer_ctx->region_type == ORT_ACC_PARALLEL) > + && (outer_ctx->region_type == ORT_ACC_PARALLEL > + || outer_ctx->region_type == ORT_ACC_SERIAL)) > { > /* Remove firstprivate and make it a copy map. */ > n->value &= ~GOVD_FIRSTPRIVATE; > @@ -6980,7 +6984,8 @@ omp_add_variable (struct gimplify_omp_ct > } > } > else if (update_data_map > - && outer_ctx->region_type == ORT_ACC_PARALLEL) > + && (outer_ctx->region_type == ORT_ACC_PARALLEL > + || outer_ctx->region_type == ORT_ACC_SERIAL)) > { > unsigned f = GOVD_MAP | GOVD_SEEN; > > @@ -7208,7 +7213,8 @@ oacc_default_clause (struct gimplify_omp > break; > > case ORT_ACC_PARALLEL: > - rkind = "parallel"; > + case ORT_ACC_SERIAL: > + rkind = ctx->region_type == ORT_ACC_PARALLEL ? "parallel" : "serial"; > > if (TREE_CODE (type) == REFERENCE_TYPE > && TREE_CODE (TREE_TYPE (type)) == POINTER_TYPE) > @@ -7828,6 +7834,7 @@ gimplify_scan_omp_clauses (tree *list_p, > case OACC_HOST_DATA: > //case OACC_PARALLEL: > //case OACC_KERNELS: > + //case OACC_SERIAL: > ctx->target_firstprivatize_array_bases = true; > default: > break; > @@ -8985,7 +8992,8 @@ gomp_needs_data_present (tree decl) > return NULL_TREE; > > if (gimplify_omp_ctxp->region_type != ORT_ACC_PARALLEL > - && gimplify_omp_ctxp->region_type != ORT_ACC_KERNELS) > + && gimplify_omp_ctxp->region_type != ORT_ACC_KERNELS > + && gimplify_omp_ctxp->region_type != ORT_ACC_SERIAL) > return NULL_TREE; > > for (ctx = gimplify_omp_ctxp->outer_context; !found_match && ctx; > @@ -9442,7 +9450,8 @@ gimplify_adjust_omp_clauses (gimple_seq > /* Data clauses associated with acc parallel 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; > @@ -9601,7 +9610,8 @@ gimplify_adjust_omp_clauses (gimple_seq > 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)) > @@ -11041,6 +11051,9 @@ gimplify_omp_workshare (tree *expr_p, gi > case OACC_PARALLEL: > ort = ORT_ACC_PARALLEL; > break; > + case OACC_SERIAL: > + ort = ORT_ACC_SERIAL; > + break; > case OACC_DATA: > ort = ORT_ACC_DATA; > break; > @@ -11115,6 +11128,10 @@ gimplify_omp_workshare (tree *expr_p, gi > 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; > @@ -12316,6 +12333,7 @@ gimplify_expr (tree *expr_p, gimple_seq > case OACC_DATA: > case OACC_KERNELS: > case OACC_PARALLEL: > + case OACC_SERIAL: > case OMP_SECTIONS: > case OMP_SINGLE: > case OMP_TARGET: > @@ -12708,6 +12726,7 @@ gimplify_expr (tree *expr_p, gimple_seq > && code != TRY_FINALLY_EXPR > && code != OACC_PARALLEL > && code != OACC_KERNELS > + && code != OACC_SERIAL > && code != OACC_DATA > && code != OACC_HOST_DATA > && code != OACC_DECLARE > Index: gcc-openacc-gcc-8-branch/gcc/omp-expand.c > =================================================================== > --- gcc-openacc-gcc-8-branch.orig/gcc/omp-expand.c > +++ gcc-openacc-gcc-8-branch/gcc/omp-expand.c > @@ -6959,6 +6959,7 @@ expand_omp_target (struct omp_region *re > switch (gimple_omp_target_kind (entry_stmt)) > { > case GF_OMP_TARGET_KIND_OACC_PARALLEL: > + case GF_OMP_TARGET_KIND_OACC_SERIAL: > oacc_parallel = true; > gcc_fallthrough (); > case GF_OMP_TARGET_KIND_REGION: > @@ -6996,16 +6997,28 @@ expand_omp_target (struct omp_region *re > entry_bb = region->entry; > exit_bb = region->exit; > > - if (gimple_omp_target_kind (entry_stmt) == GF_OMP_TARGET_KIND_OACC_KERNELS) > + switch (gimple_omp_target_kind (entry_stmt)) > { > + 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) > @@ -7214,6 +7227,7 @@ expand_omp_target (struct omp_region *re > break; > case GF_OMP_TARGET_KIND_OACC_KERNELS: > case GF_OMP_TARGET_KIND_OACC_PARALLEL: > + case GF_OMP_TARGET_KIND_OACC_SERIAL: > start_ix = BUILT_IN_GOACC_PARALLEL; > break; > case GF_OMP_TARGET_KIND_OACC_DATA: > @@ -7379,7 +7393,18 @@ expand_omp_target (struct omp_region *re > 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: > @@ -8001,6 +8026,7 @@ build_omp_regions_1 (basic_block bb, str > 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: > if (is_gimple_omp_oacc (stmt)) > @@ -8249,6 +8275,7 @@ omp_make_gimple_edges (basic_block bb, s > 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; > Index: gcc-openacc-gcc-8-branch/gcc/omp-low.c > =================================================================== > --- gcc-openacc-gcc-8-branch.orig/gcc/omp-low.c > +++ gcc-openacc-gcc-8-branch/gcc/omp-low.c > @@ -150,15 +150,17 @@ static tree scan_omp_1_op (tree *, int * > *handled_ops_p = false; \ > break; > > -/* Return true if CTX corresponds to an oacc parallel region. */ > +/* Return true if CTX corresponds to an oacc 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. */ > @@ -508,7 +510,7 @@ build_receiver_ref (tree var, bool by_re > { > tree x, field = lookup_field (var, ctx); > > - if (is_oacc_parallel (ctx)) > + if (is_oacc_parallel_or_serial (ctx)) > x = lookup_parm (var, ctx); > else > { > @@ -660,7 +662,7 @@ build_sender_ref (tree var, omp_context > static void > install_parm_decl (tree var, tree type, omp_context *ctx) > { > - if (!is_oacc_parallel (ctx)) > + if (!is_oacc_parallel_or_serial (ctx)) > return; > > splay_tree_key key = (splay_tree_key) var; > @@ -1223,7 +1225,7 @@ scan_sharing_clauses (tree clauses, omp_ > /* FIXME: The "oacc gangprivate" attribute conflicts with > the privatization of acc loops. Remove that attribute, > if present. */ > - if (!is_oacc_parallel (ctx)) > + if (!is_oacc_parallel_or_serial (ctx)) > { > tree attributes = DECL_ATTRIBUTES (new_decl); > attributes = remove_attribute ("oacc gangprivate", > @@ -1838,7 +1840,7 @@ create_omp_child_function (omp_context * > if (task_copy) > type = build_function_type_list (void_type_node, ptr_type_node, > ptr_type_node, NULL_TREE); > - else if (is_oacc_parallel (ctx)) > + else if (is_oacc_parallel_or_serial (ctx)) > { > tree *arg_types = (tree *) alloca (sizeof (tree) * map_cnt); > for (unsigned int i = 0; i < map_cnt; i++) > @@ -1918,7 +1920,7 @@ create_omp_child_function (omp_context * > DECL_CONTEXT (t) = decl; > DECL_RESULT (decl) = t; > > - if (!is_oacc_parallel (ctx)) > + if (!is_oacc_parallel_or_serial (ctx)) > { > tree data_name = get_identifier (".omp_data_i"); > t = build_decl (DECL_SOURCE_LOCATION (decl), PARM_DECL, data_name, > @@ -2409,7 +2411,7 @@ scan_omp_for (gomp_for *stmt, omp_contex > { > 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; > @@ -2638,7 +2640,7 @@ scan_omp_target (gomp_target *stmt, omp_ > bool base_pointers_restrict = false; > if (offloaded) > { > - if (!is_oacc_parallel (ctx)) > + if (!is_oacc_parallel_or_serial (ctx)) > { > create_omp_child_function (ctx, false); > gimple_omp_target_set_child_fn (stmt, ctx->cb.dst_fn); > @@ -2803,6 +2805,7 @@ check_omp_nesting_restrictions (gimple * > { > case GF_OMP_TARGET_KIND_OACC_PARALLEL: > case GF_OMP_TARGET_KIND_OACC_KERNELS: > + case GF_OMP_TARGET_KIND_OACC_SERIAL: > ok = true; > break; > > @@ -3219,6 +3222,7 @@ check_omp_nesting_restrictions (gimple * > 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: > @@ -3235,6 +3239,8 @@ check_omp_nesting_restrictions (gimple * > 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; > @@ -5263,8 +5269,10 @@ lower_oacc_reductions (location_t loc, t > 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); > @@ -6053,7 +6061,8 @@ lower_oacc_head_mark (location_t loc, tr > /* In a parallel region, loops without auto and seq clauses are > implicitly INDEPENDENT. */ > omp_context *tgt = enclosing_target_ctx (ctx); > - if ((!tgt || is_oacc_parallel (tgt)) && !(tag & (OLF_SEQ | OLF_AUTO))) > + if ((!tgt || is_oacc_parallel_or_serial (tgt)) > + && !(tag & (OLF_SEQ | OLF_AUTO))) > tag |= OLF_INDEPENDENT; > > if (tag & OLF_TILE) > @@ -8001,7 +8010,7 @@ convert_from_firstprivate_int (tree var, > static tree > append_decl_arg (tree var, tree decl_args, omp_context *ctx) > { > - if (!is_oacc_parallel (ctx)) > + if (!is_oacc_parallel_or_serial (ctx)) > return NULL_TREE; > > tree temp = lookup_parm (var, ctx); > @@ -8034,6 +8043,7 @@ lower_omp_target (gimple_stmt_iterator * > 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: > @@ -8075,7 +8085,7 @@ lower_omp_target (gimple_stmt_iterator * > > /* Determine init_cnt to finish initialize ctx. */ > > - if (is_oacc_parallel (ctx)) > + if (is_oacc_parallel_or_serial (ctx)) > { > for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c)) > switch (OMP_CLAUSE_CODE (c)) > @@ -8125,7 +8135,7 @@ lower_omp_target (gimple_stmt_iterator * > break; > > case OMP_CLAUSE_FIRSTPRIVATE: > - if (is_oacc_parallel (ctx)) > + if (is_oacc_parallel_or_serial (ctx)) > goto init_oacc_firstprivate; > init_cnt++; > break; > @@ -8326,7 +8336,7 @@ lower_omp_target (gimple_stmt_iterator * > 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); > @@ -8410,7 +8420,7 @@ lower_omp_target (gimple_stmt_iterator * > > if (offloaded) > { > - if (is_oacc_parallel (ctx)) > + if (is_oacc_parallel_or_serial (ctx)) > gcc_assert (init_cnt == map_cnt); > target_nesting_level++; > lower_omp (&tgt_body, ctx); > @@ -8744,7 +8754,7 @@ lower_omp_target (gimple_stmt_iterator * > 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)) > @@ -8849,7 +8859,7 @@ lower_omp_target (gimple_stmt_iterator * > } > > gcc_assert (map_idx == map_cnt); > - if (is_oacc_parallel (ctx)) > + if (is_oacc_parallel_or_serial (ctx)) > DECL_ARGUMENTS (child_fn) = nreverse (decl_args); > > DECL_INITIAL (TREE_VEC_ELT (t, 1)) > @@ -8889,7 +8899,7 @@ lower_omp_target (gimple_stmt_iterator * > { > t = build_fold_addr_expr_loc (loc, ctx->sender_decl); > /* fixup_child_record_type might have changed receiver_decl's type. */ > - if (!is_oacc_parallel (ctx)) > + if (!is_oacc_parallel_or_serial (ctx)) > { > t = fold_convert_loc (loc, TREE_TYPE (ctx->receiver_decl), t); > gimple_seq_add_stmt (&new_body, > @@ -9218,7 +9228,7 @@ lower_omp_target (gimple_stmt_iterator * > 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. */ > Index: gcc-openacc-gcc-8-branch/gcc/testsuite/c-c++-common/goacc/serial-dims.c > =================================================================== > --- /dev/null > +++ gcc-openacc-gcc-8-branch/gcc/testsuite/c-c++-common/goacc/serial-dims.c > @@ -0,0 +1,12 @@ > +/* Invalid use of OpenACC parallelism dimensions clauses: num_gangs, > + num_workers, vector_length with the serial construct. */ > + > +void f(void) > +{ > +#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'" } */ > + ; > +} > Index: gcc-openacc-gcc-8-branch/gcc/tree-pretty-print.c > =================================================================== > --- gcc-openacc-gcc-8-branch.orig/gcc/tree-pretty-print.c > +++ gcc-openacc-gcc-8-branch/gcc/tree-pretty-print.c > @@ -2987,6 +2987,10 @@ dump_generic_node (pretty_printer *pp, t > 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); > Index: gcc-openacc-gcc-8-branch/gcc/tree.def > =================================================================== > --- gcc-openacc-gcc-8-branch.orig/gcc/tree.def > +++ gcc-openacc-gcc-8-branch/gcc/tree.def > @@ -1096,6 +1096,12 @@ DEFTREECODE (OACC_PARALLEL, "oacc_parall > > 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. */ > Index: gcc-openacc-gcc-8-branch/libgomp/testsuite/libgomp.oacc-c-c++-common/serial-dims.c > =================================================================== > --- /dev/null > +++ gcc-openacc-gcc-8-branch/libgomp/testsuite/libgomp.oacc-c-c++-common/serial-dims.c > @@ -0,0 +1,98 @@ > +/* OpenACC dimensions with the serial construct. */ > + > +/* { dg-additional-options "-foffload-force" } */ > + > +#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 (); > +} > + > + > +int main () > +{ > + acc_init (acc_device_default); > + > + /* Serial OpenACC constructs must get launched as 1 x 1 x 1. */ > + { > + int gangs_min, gangs_max; > + int workers_min, workers_max; > + int vectors_min, vectors_max; > + int gangs_actual, workers_actual, vectors_actual; > + int i, j, k; > + > + gangs_min = workers_min = vectors_min = INT_MAX; > + gangs_max = workers_max = vectors_max = INT_MIN; > + gangs_actual = workers_actual = vectors_actual = 1; > +#pragma acc serial > + /* { dg-warning "region contains gang partitoned code but is not gang partitioned" "" { target *-*-* } 60 } */ > + /* { dg-warning "region contains worker partitoned code but is not worker partitioned" "" { target *-*-* } 60 } */ > + /* { dg-warning "region contains vector partitoned code but is not vector partitioned" "" { target *-*-* } 60 } */ > + /* { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } 60 } */ > + { > + if (acc_on_device (acc_device_nvidia)) > + { > + /* The GCC nvptx back end enforces vector_length (32). */ > + vectors_actual = 32; > + } > + else if (!acc_on_device (acc_device_host)) > + __builtin_abort (); > +#pragma acc loop gang \ > + reduction (min: gangs_min, workers_min, vectors_min) \ > + reduction (max: gangs_max, workers_max, vectors_max) > + for (i = 100 * gangs_actual; i > -100 * gangs_actual; i--) > +#pragma acc loop worker \ > + reduction (min: gangs_min, workers_min, vectors_min) \ > + reduction (max: gangs_max, workers_max, vectors_max) > + for (j = 100 * workers_actual; j > -100 * workers_actual; j--) > +#pragma acc loop vector \ > + reduction (min: gangs_min, workers_min, vectors_min) \ > + reduction (max: gangs_max, workers_max, vectors_max) > + for (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 (gangs_min != 0 || gangs_max != gangs_actual - 1 > + || workers_min != 0 || workers_max != workers_actual - 1 > + || vectors_min != 0 || vectors_max != vectors_actual - 1) > + __builtin_abort (); > + } > + } > + > + return 0; > +} >