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; +}