This is the mail archive of the gcc-patches@gcc.gnu.org mailing list for the GCC project.


Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]
Other format: [Raw text]

[PATCH, og8] Add OpenACC 2.6 `serial' construct support


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


Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]