| Message ID | alpine.DEB.2.21.9999.1812161533170.99920@build7-trusty-cs.sje.mentorg.com |
|---|---|
| State | Accepted |
| Headers | show |
| Series | [og8] Add OpenACC 2.6 `serial' construct support | expand |
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. > <https://gcc.gnu.org/ml/gcc-patches/2015-11/msg00245.html>). 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 <limits.h> > +#include <openacc.h> > +#include <gomp-constants.h> > + > +/* 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; > +} >
Hi Chung-Lin, > 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... Thank you for your review. I think your input is valuable and it may be worth it to revise the implementation at the time of merging it to trunk. However I think that perfect is the enemy of good, so for the time being I have committed this change. Maciej
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 <limits.h> +#include <openacc.h> +#include <gomp-constants.h> + +/* 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; +}
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. <https://gcc.gnu.org/ml/gcc-patches/2015-11/msg00245.html>). 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