Add OpenACC 2.6 `serial' construct support

    The `serial' construct (cf. section 2.5.3 of the OpenACC 2.6 standard)
    is equivalent to a `parallel' construct with clauses `num_gangs(1)
    num_workers(1) vector_length(1)' implied.
    These clauses are therefore not supported with the `serial'
    construct. All the remaining clauses accepted with `parallel' are also
    accepted with `serial'.

    The `serial' construct is implemented like `parallel', except for
    hardcoding dimensions rather than taking them from the relevant
    clauses, in `expand_omp_target'.

    Separate codes are used to denote the `serial' construct throughout the
    middle end, even though the mapping of `serial' to an equivalent
    `parallel' construct could have been done in the individual language
    frontends. In particular, this allows to distinguish between compute
    constructs in warnings, error messages, dumps etc.

    2019-11-12  Maciej W. Rozycki  <macro@codesourcery.com>
                Tobias Burnus  <tobias@codesourcery.com>
                Frederik Harwath  <frederik@codesourcery.com>
                Thomas Schwinge  <thomas@codesourcery.com>

    gcc/
    * gimple.h (gf_mask): Add GF_OMP_TARGET_KIND_OACC_SERIAL
    enumeration constant.
    (is_gimple_omp_oacc): Handle GF_OMP_TARGET_KIND_OACC_SERIAL.
    (is_gimple_omp_offloaded): Likewise.
    * gimplify.c (omp_region_type): Add ORT_ACC_SERIAL enumeration
    constant.  Adjust the value of ORT_NONE accordingly.
    (is_gimple_stmt): Handle OACC_SERIAL.
    (oacc_default_clause): Handle ORT_ACC_SERIAL.
    (gomp_needs_data_present): Likewise.
    (gimplify_adjust_omp_clauses): Likewise.
    (gimplify_omp_workshare): Handle OACC_SERIAL.
    (gimplify_expr): Likewise.
    * omp-expand.c (expand_omp_target):
    Handle GF_OMP_TARGET_KIND_OACC_SERIAL.
    (build_omp_regions_1, omp_make_gimple_edges): Likewise.
    * omp-low.c (is_oacc_parallel): Rename function to...
    (is_oacc_parallel_or_serial): ... this.
    Handle GF_OMP_TARGET_KIND_OACC_SERIAL.
    (scan_sharing_clauses): Adjust accordingly.
    (scan_omp_for): Likewise.
    (lower_oacc_head_mark): Likewise.
    (convert_from_firstprivate_int): Likewise.
    (lower_omp_target): Likewise.
    (check_omp_nesting_restrictions): Handle
    GF_OMP_TARGET_KIND_OACC_SERIAL.
    (lower_oacc_reductions): Likewise.
    (lower_omp_target): Likewise.
    * tree.def (OACC_SERIAL): New tree code.
    * tree-pretty-print.c (dump_generic_node): Handle OACC_SERIAL.

    * doc/generic.texi (OpenACC): Document OACC_SERIAL.

    gcc/c-family/
    * c-pragma.h (pragma_kind): Add PRAGMA_OACC_SERIAL enumeration
    constant.
    * c-pragma.c (oacc_pragmas): Add "serial" entry.

    gcc/c/
    * c-parser.c (OACC_SERIAL_CLAUSE_MASK): New macro.
    (c_parser_oacc_kernels_parallel): Rename function to...
    (c_parser_oacc_compute): ... this.  Handle PRAGMA_OACC_SERIAL.
    (c_parser_omp_construct): Update accordingly.

    gcc/cp/
    * constexpr.c (potential_constant_expression_1): Handle
    OACC_SERIAL.
    * parser.c (OACC_SERIAL_CLAUSE_MASK): New macro.
    (cp_parser_oacc_kernels_parallel): Rename function to...
    (cp_parser_oacc_compute): ... this.  Handle PRAGMA_OACC_SERIAL.
    (cp_parser_omp_construct): Update accordingly.
    (cp_parser_pragma): Handle PRAGMA_OACC_SERIAL.  Fix alphabetic
    order.
    * pt.c (tsubst_expr): Handle OACC_SERIAL.

    gcc/fortran/
    * gfortran.h (gfc_statement): Add ST_OACC_SERIAL_LOOP,
    ST_OACC_END_SERIAL_LOOP, ST_OACC_SERIAL and ST_OACC_END_SERIAL
    enumeration constants.
    (gfc_exec_op): Add EXEC_OACC_SERIAL_LOOP and EXEC_OACC_SERIAL
    enumeration constants.
    * match.h (gfc_match_oacc_serial): New prototype.
    (gfc_match_oacc_serial_loop): Likewise.
    * dump-parse-tree.c (show_omp_node, show_code_node): Handle
    EXEC_OACC_SERIAL_LOOP and EXEC_OACC_SERIAL.
    * match.c (match_exit_cycle): Handle EXEC_OACC_SERIAL_LOOP.
    * openmp.c (OACC_SERIAL_CLAUSES): New macro.
    (gfc_match_oacc_serial_loop): New function.
    (gfc_match_oacc_serial): Likewise.
    (oacc_is_loop): Handle EXEC_OACC_SERIAL_LOOP.
    (resolve_omp_clauses): Handle EXEC_OACC_SERIAL.
    (oacc_code_to_statement): Handle EXEC_OACC_SERIAL and
    EXEC_OACC_SERIAL_LOOP.
    (gfc_resolve_oacc_directive): Likewise.
    * parse.c (decode_oacc_directive) <'s'>: Add case for "serial"
    and "serial loop".
    (next_statement): Handle ST_OACC_SERIAL_LOOP and ST_OACC_SERIAL.
    (gfc_ascii_statement): Likewise.  Handle ST_OACC_END_SERIAL_LOOP
    and ST_OACC_END_SERIAL.
    (parse_oacc_structured_block): Handle ST_OACC_SERIAL.
    (parse_oacc_loop): Handle ST_OACC_SERIAL_LOOP and
    ST_OACC_END_SERIAL_LOOP.
    (parse_executable): Handle ST_OACC_SERIAL_LOOP and
    ST_OACC_SERIAL.
    (is_oacc): Handle EXEC_OACC_SERIAL_LOOP and EXEC_OACC_SERIAL.
    * resolve.c (gfc_resolve_blocks, gfc_resolve_code): Likewise.
    * st.c (gfc_free_statement): Likewise.
    * trans-openmp.c (gfc_trans_oacc_construct): Handle
    EXEC_OACC_SERIAL.
    (gfc_trans_oacc_combined_directive): Handle
    EXEC_OACC_SERIAL_LOOP.
    (gfc_trans_oacc_directive): Handle EXEC_OACC_SERIAL_LOOP and
    EXEC_OACC_SERIAL.
    * trans.c (trans_code): Likewise.

    gcc/testsuite/
    * c-c++-common/goacc/parallel-dims.c: New test.
    * gfortran.dg/goacc/parallel-dims.f90: New test.

    libgomp/
    * testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: New test.
    * testsuite/libgomp.oacc-fortran/parallel-dims-aux.c: New test.
    * testsuite/libgomp.oacc-fortran/parallel-dims.f89: New test.
    * testsuite/libgomp.oacc-fortran/parallel-dims-2.f90: New test.

Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>


Co-Authored-By: Frederik Harwath <frederik@codesourcery.com>
Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com>
Co-Authored-By: Tobias Burnus <tobias@codesourcery.com>

From-SVN: r278082

1 files changed, 5 insertions, 0 deletions

diff --git a/gcc/doc/generic.texi b/gcc/doc/generic.texi
index 94e339c..badaaec 100644
--- a/gcc/doc/generic.texi
+++ b/gcc/doc/generic.texi
@@ -2388,6 +2388,7 @@ compilation.
 @tindex OACC_KERNELS
 @tindex OACC_LOOP
 @tindex OACC_PARALLEL
+@tindex OACC_SERIAL
 @tindex OACC_UPDATE
 
 All the statements starting with @code{OACC_} represent directives and
@@ -2432,6 +2433,10 @@ See the description of the @code{OMP_FOR} code.
 
 Represents @code{#pragma acc parallel [clause1 @dots{} clauseN]}.
 
+@item OACC_SERIAL
+
+Represents @code{#pragma acc serial [clause1 @dots{} clauseN]}.
+
 @item OACC_UPDATE
 
 Represents @code{#pragma acc update [clause1 @dots{} clauseN]}.