gimplify.c (gimplify_scan_omp_clauses): Add support for OMP_CLAUSE_TILE.

	gcc/
	* gimplify.c (gimplify_scan_omp_clauses): Add support for
	OMP_CLAUSE_TILE.  Update handling of OMP_CLAUSE_INDEPENDENT.
	(gimplify_adjust_omp_clauses): Likewise.
	* omp-low.c (scan_sharing_clauses): Add support for OMP_CLAUSE_TILE.
	* tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_TILE.
	* tree-pretty-print.c (dump_omp_clause): Handle OMP_CLAUSE_TILE.
	* tree.c (omp_clause_num_ops): Add an entry for OMP_CLAUSE_TILE.
	(omp_clause_code_name): Likewise.
	(walk_tree_1): Handle OMP_CLAUSE_TILE.
	* tree.h (OMP_TILE_LIST): New macro.

	gcc/c-family/
	* c-omp.c (c_oacc_split_loop_clauses): Make TILE, GANG, WORKER, VECTOR,
	AUTO, SEQ, INDEPENDENT and PRIVATE loop clauses.  Associate REDUCTION
	clauses with parallel and kernels and loops.
	* c-pragma.h (enum pragma_omp_clause): Add entries for
	PRAGMA_OACC_CLAUSE_{INDEPENDENT,TILE,DEFAULT}.
	* pt.c (tsubst_omp_clauses): Add support for OMP_CLAUSE_{NUM_GANGS,
	NUM_WORKERS,VECTOR_LENGTH,GANG,WORKER,VECTOR,ASYNC,WAIT,TILE,AUTO,
	INDEPENDENT,SEQ}. 
	(tsubst_expr): Add support for OMP_CLAUSE_{KERNELS,PARALLEL,LOOP}.

	gcc/c/
	* c-parser.c (c_parser_omp_clause_name): Add support for
	PRAGMA_OACC_CLAUSE_INDEPENDENT and PRAGMA_OACC_CLAUSE_TILE.
	(c_parser_omp_clause_default): Add is_oacc argument. Handle
	default(none) in OpenACC.
	(c_parser_oacc_shape_clause): Allow pointer variables as gang static
	arguments.
	(c_parser_oacc_clause_tile): New function.
	(c_parser_oacc_all_clauses): Add support for OMP_CLAUSE_DEFAULT,
	OMP_CLAUSE_INDEPENDENT and OMP_CLAUSE_TILE.
	(OACC_LOOP_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_{PRIVATE,INDEPENDENT,
	TILE}.
	(OACC_KERNELS_MASK): Add PRAGMA_OACC_CLAUSE_DEFAULT.
	(OACC_PARALLEL_MASK): Add PRAGMA_OACC_CLAUSE_{DEFAULT,PRIVATE,
	FIRSTPRIVATE}.
	(c_parser_omp_all_clauses): Update call to c_parser_omp_clause_default.
	(c_parser_oacc_update): Update the error message for missing clauses.
	* c-typeck.c (c_finish_omp_clauses): Add support for OMP_CLAUSE_TILE
	and OMP_CLAUSE_INDEPENDENT.

	gcc/cp/
	* parser.c (cp_parser_omp_clause_name): Add support for
	PRAGMA_OACC_CLAUSE_INDEPENDENT and PRAGMA_OACC_CLAUSE_TILE.
	(cp_parser_oacc_shape_clause): Allow pointer variables as gang static
	arguments.
	(cp_parser_oacc_clause_tile): New function.
	(cp_parser_omp_clause_default): Add is_oacc argument. Handle
	default(none) in OpenACC.
	(cp_parser_oacc_all_clauses): Add support for
	(cp_parser_omp_all_clauses): Update call to
	cp_parser_omp_clause_default.
	PRAGMA_OACC_CLAUSE_{DEFAULT,INDEPENDENT,TILE,PRIVATE,FIRSTPRIVATE}.
	(OACC_LOOP_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_{PRIVATE,INDEPENDENT,
	TILE}.
	(OACC_KERNELS_MASK): Add PRAGMA_OACC_CLAUSE_DEFAULT.
	(OACC_PARALLEL_MASK): Add PRAGMA_OACC_CLAUSE_{DEFAULT,PRIVATE,
	FIRSTPRIVATE}.
	(cp_parser_oacc_update): Update the error message for missing clauses.
	* semantics.c (finish_omp_clauses): Add support for
	OMP_CLAUSE_INDEPENDENT and OMP_CLAUSE_TILE.

	gcc/fortran/
	* openmp.c (gfc_match_omp_clauses): Update support for the tile
	and default clauses in OpenACC.
	(gfc_match_oacc_update): Error when data clauses are supplied.
	(oacc_compatible_clauses): Delete.
	(resolve_omp_clauses): Give special care for OpenACC reductions.
	Also update error reporting for the tile clause.
	(resolve_oacc_loop_blocks): Update error reporting for the tile clause.
	* trans-openmp.c (gfc_trans_omp_clauses): Update OMP_CLAUSE_SEQ. Add
	OMP_CLAUSE_{AUTO,TILE} and add support the the gang static argument.
	(gfc_trans_oacc_combined_directive): Update the list of clauses which
	are split to acc loops.

	gcc/testsuite/
	* c-c++-common/goacc/combined-directives.c: New test.
	* c-c++-common/goacc/loop-clauses.c: New test.
	* c-c++-common/goacc/tile.c: New test.
	* c-c++-common/goacc/loop-shape.c: Add test for pointer variable
	as gang static arguments.
	* c-c++-common/goacc/update-1.c: Adjust expected error message.
	* g++.dg/goacc/template.C: New test.
	* gfortran.dg/goacc/combined-directives.f90: New test.
	* gfortran.dg/goacc/default.f95: New test.
	* gfortran.dg/goacc/default_none.f95: New test.
	* gfortran.dg/goacc/firstprivate-1.f95: New test.
	* gfortran.dg/goacc/gang-static.f95: New test.
	* gfortran.dg/goacc/kernels-loop-inner.f95: New test.
	* gfortran.dg/goacc/kernels-loops-adjacent.f95: New test.
	* gfortran.dg/goacc/list.f95: Update test.
	* gfortran.dg/goacc/loop-2.f95: Likewise.
	* gfortran.dg/goacc/loop-4.f95: New test.
	* gfortran.dg/goacc/loop-5.f95: New test.
	* gfortran.dg/goacc/loop-6.f95: New test.
	* gfortran.dg/goacc/loop-tree-1.f90: Update test.
	* gfortran.dg/goacc/multi-clause.f90: New test.
	* gfortran.dg/goacc/parallel-tree.f95: Update test.
	* gfortran.dg/goacc/update.f95: New test.


Co-Authored-By: James Norris <jnorris@codesourcery.com>
Co-Authored-By: Nathan Sidwell <nathan@codesourcery.com>
Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com>
Co-Authored-By: Tom de Vries <tom@codesourcery.com>

From-SVN: r229832

43 files changed, 1840 insertions, 98 deletions

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index 94e597e..f01e8ec 100644
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,19 @@
+2015-11-05  Cesar Philippidis  <cesar@codesourcery.com>
+	    Thomas Schwinge  <thomas@codesourcery.com>
+	    James Norris  <jnorris@codesourcery.com>
+
+
+	* gimplify.c (gimplify_scan_omp_clauses): Add support for
+	OMP_CLAUSE_TILE.  Update handling of OMP_CLAUSE_INDEPENDENT.
+	(gimplify_adjust_omp_clauses): Likewise.
+	* omp-low.c (scan_sharing_clauses): Add support for OMP_CLAUSE_TILE.
+	* tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_TILE.
+	* tree-pretty-print.c (dump_omp_clause): Handle OMP_CLAUSE_TILE.
+	* tree.c (omp_clause_num_ops): Add an entry for OMP_CLAUSE_TILE.
+	(omp_clause_code_name): Likewise.
+	(walk_tree_1): Handle OMP_CLAUSE_TILE.
+	* tree.h (OMP_TILE_LIST): New macro.
+
 2015-11-05  Martin Sebor  <msebor@redhat.com>
 
 	PR c++/67942
diff --git a/gcc/c-family/ChangeLog b/gcc/c-family/ChangeLog
index c21f504..cc9a642 100644
--- a/gcc/c-family/ChangeLog
+++ b/gcc/c-family/ChangeLog
@@ -1,3 +1,17 @@
+2015-11-05  Cesar Philippidis  <cesar@codesourcery.com>
+	    Thomas Schwinge  <thomas@codesourcery.com>
+	    James Norris  <jnorris@codesourcery.com>
+
+	* c-omp.c (c_oacc_split_loop_clauses): Make TILE, GANG, WORKER, VECTOR,
+	AUTO, SEQ, INDEPENDENT and PRIVATE loop clauses.  Associate REDUCTION
+	clauses with parallel and kernels and loops.
+	* c-pragma.h (enum pragma_omp_clause): Add entries for
+	PRAGMA_OACC_CLAUSE_{INDEPENDENT,TILE,DEFAULT}.
+	* pt.c (tsubst_omp_clauses): Add support for OMP_CLAUSE_{NUM_GANGS,
+	NUM_WORKERS,VECTOR_LENGTH,GANG,WORKER,VECTOR,ASYNC,WAIT,TILE,AUTO,
+	INDEPENDENT,SEQ}.
+	(tsubst_expr): Add support for OMP_CLAUSE_{KERNELS,PARALLEL,LOOP}.
+
 2015-11-05  Martin Sebor  <msebor@redhat.com>
 
 	PR c++/67942
diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c
index a51611c..3e93b59 100644
--- a/gcc/c-family/c-omp.c
+++ b/gcc/c-family/c-omp.c
@@ -867,7 +867,7 @@ c_omp_check_loop_iv_exprs (location_t stmt_loc, tree declv, tree decl,
 tree
 c_oacc_split_loop_clauses (tree clauses, tree *not_loop_clauses)
 {
-  tree next, loop_clauses;
+  tree next, loop_clauses, t;
 
   loop_clauses = *not_loop_clauses = NULL_TREE;
   for (; clauses ; clauses = next)
@@ -876,12 +876,27 @@ c_oacc_split_loop_clauses (tree clauses, tree *not_loop_clauses)
 
       switch (OMP_CLAUSE_CODE (clauses))
         {
+	  /* Loop clauses.  */
 	case OMP_CLAUSE_COLLAPSE:
-	case OMP_CLAUSE_REDUCTION:
+	case OMP_CLAUSE_TILE:
+	case OMP_CLAUSE_GANG:
+	case OMP_CLAUSE_WORKER:
+	case OMP_CLAUSE_VECTOR:
+	case OMP_CLAUSE_AUTO:
+	case OMP_CLAUSE_SEQ:
+	case OMP_CLAUSE_INDEPENDENT:
+	case OMP_CLAUSE_PRIVATE:
 	  OMP_CLAUSE_CHAIN (clauses) = loop_clauses;
 	  loop_clauses = clauses;
 	  break;
 
+	  /* Reductions belong in both constructs.  */
+	case OMP_CLAUSE_REDUCTION:
+	  t = copy_node (clauses);
+	  OMP_CLAUSE_CHAIN (t) = loop_clauses;
+	  loop_clauses = t;
+
+	  /* Parallel/kernels clauses.  */
 	default:
 	  OMP_CLAUSE_CHAIN (clauses) = *not_loop_clauses;
 	  *not_loop_clauses = clauses;
diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h
index 69e7392..953c4e3 100644
--- a/gcc/c-family/c-pragma.h
+++ b/gcc/c-family/c-pragma.h
@@ -153,6 +153,7 @@ enum pragma_omp_clause {
   PRAGMA_OACC_CLAUSE_DEVICEPTR,
   PRAGMA_OACC_CLAUSE_GANG,
   PRAGMA_OACC_CLAUSE_HOST,
+  PRAGMA_OACC_CLAUSE_INDEPENDENT,
   PRAGMA_OACC_CLAUSE_NUM_GANGS,
   PRAGMA_OACC_CLAUSE_NUM_WORKERS,
   PRAGMA_OACC_CLAUSE_PRESENT,
@@ -162,6 +163,7 @@ enum pragma_omp_clause {
   PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE,
   PRAGMA_OACC_CLAUSE_SELF,
   PRAGMA_OACC_CLAUSE_SEQ,
+  PRAGMA_OACC_CLAUSE_TILE,
   PRAGMA_OACC_CLAUSE_VECTOR,
   PRAGMA_OACC_CLAUSE_VECTOR_LENGTH,
   PRAGMA_OACC_CLAUSE_WAIT,
@@ -169,6 +171,7 @@ enum pragma_omp_clause {
   PRAGMA_OACC_CLAUSE_COLLAPSE = PRAGMA_OMP_CLAUSE_COLLAPSE,
   PRAGMA_OACC_CLAUSE_COPYIN = PRAGMA_OMP_CLAUSE_COPYIN,
   PRAGMA_OACC_CLAUSE_DEVICE = PRAGMA_OMP_CLAUSE_DEVICE,
+  PRAGMA_OACC_CLAUSE_DEFAULT = PRAGMA_OMP_CLAUSE_DEFAULT,
   PRAGMA_OACC_CLAUSE_FIRSTPRIVATE = PRAGMA_OMP_CLAUSE_FIRSTPRIVATE,
   PRAGMA_OACC_CLAUSE_IF = PRAGMA_OMP_CLAUSE_IF,
   PRAGMA_OACC_CLAUSE_PRIVATE = PRAGMA_OMP_CLAUSE_PRIVATE,
diff --git a/gcc/c/ChangeLog b/gcc/c/ChangeLog
index 7839024..c734a81 100644
--- a/gcc/c/ChangeLog
+++ b/gcc/c/ChangeLog
@@ -1,3 +1,26 @@
+2015-11-05  Cesar Philippidis  <cesar@codesourcery.com>
+	    Thomas Schwinge  <thomas@codesourcery.com>
+	    James Norris  <jnorris@codesourcery.com>
+
+	* c-parser.c (c_parser_omp_clause_name): Add support for
+	PRAGMA_OACC_CLAUSE_INDEPENDENT and PRAGMA_OACC_CLAUSE_TILE.
+	(c_parser_omp_clause_default): Add is_oacc argument. Handle
+	default(none) in OpenACC.
+	(c_parser_oacc_shape_clause): Allow pointer variables as gang static
+	arguments.
+	(c_parser_oacc_clause_tile): New function.
+	(c_parser_oacc_all_clauses): Add support for OMP_CLAUSE_DEFAULT,
+	OMP_CLAUSE_INDEPENDENT and OMP_CLAUSE_TILE.
+	(OACC_LOOP_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_{PRIVATE,INDEPENDENT,
+	TILE}.
+	(OACC_KERNELS_MASK): Add PRAGMA_OACC_CLAUSE_DEFAULT.
+	(OACC_PARALLEL_MASK): Add PRAGMA_OACC_CLAUSE_{DEFAULT,PRIVATE,
+	FIRSTPRIVATE}.
+	(c_parser_omp_all_clauses): Update call to c_parser_omp_clause_default.
+	(c_parser_oacc_update): Update the error message for missing clauses.
+	* c-typeck.c (c_finish_omp_clauses): Add support for OMP_CLAUSE_TILE
+	and OMP_CLAUSE_INDEPENDENT.
+
 2015-11-05  Marek Polacek  <polacek@redhat.com>
 
 	PR c/68090
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index e2f55b3..23d0107 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -10008,6 +10008,8 @@ c_parser_omp_clause_name (c_parser *parser)
 	case 'i':
 	  if (!strcmp ("inbranch", p))
 	    result = PRAGMA_OMP_CLAUSE_INBRANCH;
+	  else if (!strcmp ("independent", p))
+	    result = PRAGMA_OACC_CLAUSE_INDEPENDENT;
 	  else if (!strcmp ("is_device_ptr", p))
 	    result = PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR;
 	  break;
@@ -10104,6 +10106,8 @@ c_parser_omp_clause_name (c_parser *parser)
 	    result = PRAGMA_OMP_CLAUSE_THREAD_LIMIT;
 	  else if (!strcmp ("threads", p))
 	    result = PRAGMA_OMP_CLAUSE_THREADS;
+	  else if (!strcmp ("tile", p))
+	    result = PRAGMA_OACC_CLAUSE_TILE;
 	  else if (!strcmp ("to", p))
 	    result = PRAGMA_OMP_CLAUSE_TO;
 	  break;
@@ -10541,10 +10545,13 @@ c_parser_omp_clause_copyprivate (c_parser *parser, tree list)
 }
 
 /* OpenMP 2.5:
-   default ( shared | none ) */
+   default ( shared | none )
+
+   OpenACC 2.0:
+   default (none) */
 
 static tree
-c_parser_omp_clause_default (c_parser *parser, tree list)
+c_parser_omp_clause_default (c_parser *parser, tree list, bool is_oacc)
 {
   enum omp_clause_default_kind kind = OMP_CLAUSE_DEFAULT_UNSPECIFIED;
   location_t loc = c_parser_peek_token (parser)->location;
@@ -10565,7 +10572,7 @@ c_parser_omp_clause_default (c_parser *parser, tree list)
 	  break;
 
 	case 's':
-	  if (strcmp ("shared", p) != 0)
+	  if (strcmp ("shared", p) != 0 || is_oacc)
 	    goto invalid_kind;
 	  kind = OMP_CLAUSE_DEFAULT_SHARED;
 	  break;
@@ -10579,7 +10586,10 @@ c_parser_omp_clause_default (c_parser *parser, tree list)
   else
     {
     invalid_kind:
-      c_parser_error (parser, "expected %<none%> or %<shared%>");
+      if (is_oacc)
+	c_parser_error (parser, "expected %<none%>");
+      else
+	c_parser_error (parser, "expected %<none%> or %<shared%>");
     }
   c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, "expected %<)%>");
 
@@ -11236,7 +11246,10 @@ c_parser_oacc_shape_clause (c_parser *parser, omp_clause_code kind,
 		}
 
 	      /* Check for the '*' argument.  */
-	      if (c_parser_next_token_is (parser, CPP_MULT))
+	      if (c_parser_next_token_is (parser, CPP_MULT)
+		  && (c_parser_peek_2nd_token (parser)->type == CPP_COMMA
+		      || c_parser_peek_2nd_token (parser)->type
+		         == CPP_CLOSE_PAREN))
 		{
 		  c_parser_consume_token (parser);
 		  ops[idx] = integer_minus_one_node;
@@ -11378,6 +11391,79 @@ c_parser_oacc_clause_async (c_parser *parser, tree list)
   return list;
 }
 
+/* OpenACC 2.0:
+   tile ( size-expr-list ) */
+
+static tree
+c_parser_oacc_clause_tile (c_parser *parser, tree list)
+{
+  tree c, expr = error_mark_node;
+  location_t loc, expr_loc;
+  tree tile = NULL_TREE;
+
+  check_no_duplicate_clause (list, OMP_CLAUSE_TILE, "tile");
+
+  loc = c_parser_peek_token (parser)->location;
+  if (!c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>"))
+    return list;
+
+  do
+    {
+      if (c_parser_next_token_is (parser, CPP_MULT)
+	  && (c_parser_peek_2nd_token (parser)->type == CPP_COMMA
+	      || c_parser_peek_2nd_token (parser)->type == CPP_CLOSE_PAREN))
+	{
+	  c_parser_consume_token (parser);
+	  expr = integer_minus_one_node;
+	}
+      else
+	{
+	  expr_loc = c_parser_peek_token (parser)->location;
+	  expr = c_parser_expr_no_commas (parser, NULL).value;
+
+	  if (expr == error_mark_node)
+	    {
+	      c_parser_skip_until_found (parser, CPP_CLOSE_PAREN,
+					 "expected %<)%>");
+	      return list;
+	    }
+
+	  if (!INTEGRAL_TYPE_P (TREE_TYPE (expr)))
+	    {
+	      c_parser_error (parser, "%<tile%> value must be integral");
+	      return list;
+	    }
+
+	  mark_exp_read (expr);
+	  expr = c_fully_fold (expr, false, NULL);
+
+	  /* Attempt to statically determine when expr isn't positive.  */
+	  c = fold_build2_loc (expr_loc, LE_EXPR, boolean_type_node, expr,
+			       build_int_cst (TREE_TYPE (expr), 0));
+	  protected_set_expr_location (c, expr_loc);
+	  if (c == boolean_true_node)
+	    {
+	      warning_at (expr_loc, 0,"%<tile%> value must be positive");
+	      expr = integer_one_node;
+	    }
+	}
+
+      tile = tree_cons (NULL_TREE, expr, tile);
+      if (c_parser_next_token_is (parser, CPP_COMMA))
+	c_parser_consume_token (parser);
+    }
+  while (c_parser_next_token_is_not (parser, CPP_CLOSE_PAREN));
+
+  /* Consume the trailing ')'.  */
+  c_parser_consume_token (parser);
+
+  c = build_omp_clause (loc, OMP_CLAUSE_TILE);
+  tile = nreverse (tile);
+  OMP_CLAUSE_TILE_LIST (c) = tile;
+  OMP_CLAUSE_CHAIN (c) = list;
+  return c;
+}
+
 /* OpenACC:
    wait ( int-expr-list ) */
 
@@ -12605,6 +12691,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
 	  clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
 	  c_name = "delete";
 	  break;
+	case PRAGMA_OMP_CLAUSE_DEFAULT:
+	  clauses = c_parser_omp_clause_default (parser, clauses, true);
+	  c_name = "default";
+	  break;
 	case PRAGMA_OACC_CLAUSE_DEVICE:
 	  clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
 	  c_name = "device";
@@ -12630,6 +12720,11 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
 	  clauses = c_parser_omp_clause_if (parser, clauses, false);
 	  c_name = "if";
 	  break;
+	case PRAGMA_OACC_CLAUSE_INDEPENDENT:
+	  clauses = c_parser_oacc_simple_clause (parser, OMP_CLAUSE_INDEPENDENT,
+						clauses);
+	  c_name = "independent";
+	  break;
 	case PRAGMA_OACC_CLAUSE_NUM_GANGS:
 	  clauses = c_parser_omp_clause_num_gangs (parser, clauses);
 	  c_name = "num_gangs";
@@ -12675,6 +12770,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
 						clauses);
 	  c_name = "seq";
 	  break;
+	case PRAGMA_OACC_CLAUSE_TILE:
+	  clauses = c_parser_oacc_clause_tile (parser, clauses);
+	  c_name = "tile";
+	  break;
 	case PRAGMA_OACC_CLAUSE_VECTOR:
 	  c_name = "vector";
 	  clauses = c_parser_oacc_shape_clause (parser, OMP_CLAUSE_VECTOR,
@@ -12756,7 +12855,7 @@ c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask,
 	  c_name = "copyprivate";
 	  break;
 	case PRAGMA_OMP_CLAUSE_DEFAULT:
-	  clauses = c_parser_omp_clause_default (parser, clauses);
+	  clauses = c_parser_omp_clause_default (parser, clauses, false);
 	  c_name = "default";
 	  break;
 	case PRAGMA_OMP_CLAUSE_FIRSTPRIVATE:
@@ -13169,13 +13268,15 @@ c_parser_oacc_enter_exit_data (c_parser *parser, bool enter)
 
 #define OACC_LOOP_CLAUSE_MASK						\
 	( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COLLAPSE)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRIVATE)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_GANG)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WORKER)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_AUTO)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_INDEPENDENT) 	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ)			\
-	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) )
-
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_TILE) )
 static tree
 c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
 		    omp_clause_mask mask, tree *cclauses)
@@ -13220,6 +13321,7 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
 	| (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_DEVICEPTR)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF)			\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT)		\
@@ -13235,8 +13337,11 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
 	| (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_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_NUM_GANGS)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT)		\
@@ -13318,7 +13423,7 @@ c_parser_oacc_update (c_parser *parser)
     {
       error_at (loc,
 		"%<#pragma acc update%> must contain at least one "
-		"%<device%> or %<host/self%> clause");
+		"%<device%> or %<host%> or %<self%> clause");
       return;
     }
 
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index ba1a8d8..33f5ead 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -13011,10 +13011,12 @@ c_finish_omp_clauses (tree clauses, bool is_omp, bool declare_simd)
 	case OMP_CLAUSE_ASYNC:
 	case OMP_CLAUSE_WAIT:
 	case OMP_CLAUSE_AUTO:
+	case OMP_CLAUSE_INDEPENDENT:
 	case OMP_CLAUSE_SEQ:
 	case OMP_CLAUSE_GANG:
 	case OMP_CLAUSE_WORKER:
 	case OMP_CLAUSE_VECTOR:
+	case OMP_CLAUSE_TILE:
 	  pc = &OMP_CLAUSE_CHAIN (c);
 	  continue;
 
diff --git a/gcc/cp/ChangeLog b/gcc/cp/ChangeLog
index 1741fa2..e780c86 100644
--- a/gcc/cp/ChangeLog
+++ b/gcc/cp/ChangeLog
@@ -1,3 +1,27 @@
+2015-11-05  Cesar Philippidis  <cesar@codesourcery.com>
+	    Thomas Schwinge  <thomas@codesourcery.com>
+	    James Norris  <jnorris@codesourcery.com>
+
+	* parser.c (cp_parser_omp_clause_name): Add support for
+	PRAGMA_OACC_CLAUSE_INDEPENDENT and PRAGMA_OACC_CLAUSE_TILE.
+	(cp_parser_oacc_shape_clause): Allow pointer variables as gang static
+	arguments.
+	(cp_parser_oacc_clause_tile): New function.
+	(cp_parser_omp_clause_default): Add is_oacc argument. Handle
+	default(none) in OpenACC.
+	(cp_parser_oacc_all_clauses): Add support for
+	(cp_parser_omp_all_clauses): Update call to
+	cp_parser_omp_clause_default.
+	PRAGMA_OACC_CLAUSE_{DEFAULT,INDEPENDENT,TILE,PRIVATE,FIRSTPRIVATE}.
+	(OACC_LOOP_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_{PRIVATE,INDEPENDENT,
+	TILE}.
+	(OACC_KERNELS_MASK): Add PRAGMA_OACC_CLAUSE_DEFAULT.
+	(OACC_PARALLEL_MASK): Add PRAGMA_OACC_CLAUSE_{DEFAULT,PRIVATE,
+	FIRSTPRIVATE}.
+	(cp_parser_oacc_update): Update the error message for missing clauses.
+	* semantics.c (finish_omp_clauses): Add support for
+	OMP_CLAUSE_INDEPENDENT and OMP_CLAUSE_TILE.
+
 2015-11-05  Martin Sebor  <msebor@redhat.com>
 
 	PR c++/67942
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index d439c06..c6f5729 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -29125,6 +29125,8 @@ cp_parser_omp_clause_name (cp_parser *parser)
 	case 'i':
 	  if (!strcmp ("inbranch", p))
 	    result = PRAGMA_OMP_CLAUSE_INBRANCH;
+	  else if (!strcmp ("independent", p))
+	    result = PRAGMA_OACC_CLAUSE_INDEPENDENT;
 	  else if (!strcmp ("is_device_ptr", p))
 	    result = PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR;
 	  break;
@@ -29219,6 +29221,8 @@ cp_parser_omp_clause_name (cp_parser *parser)
 	    result = PRAGMA_OMP_CLAUSE_THREAD_LIMIT;
 	  else if (!strcmp ("threads", p))
 	    result = PRAGMA_OMP_CLAUSE_THREADS;
+	  else if (!strcmp ("tile", p))
+	    result = PRAGMA_OACC_CLAUSE_TILE;
 	  else if (!strcmp ("to", p))
 	    result = PRAGMA_OMP_CLAUSE_TO;
 	  break;
@@ -29684,7 +29688,10 @@ cp_parser_oacc_shape_clause (cp_parser *parser, omp_clause_code kind,
 		}
 
 	      /* Check for the '*' argument.  */
-	      if (cp_lexer_next_token_is (lexer, CPP_MULT))
+	      if (cp_lexer_next_token_is (lexer, CPP_MULT)
+		  && (cp_lexer_nth_token_is (parser->lexer, 2, CPP_COMMA)
+		      || cp_lexer_nth_token_is (parser->lexer, 2,
+						CPP_CLOSE_PAREN)))
 		{
 		  cp_lexer_consume_token (lexer);
 		  ops[idx] = integer_minus_one_node;
@@ -29752,6 +29759,52 @@ cp_parser_oacc_shape_clause (cp_parser *parser, omp_clause_code kind,
   return list;
 }
 
+/* OpenACC 2.0:
+   tile ( size-expr-list ) */
+
+static tree
+cp_parser_oacc_clause_tile (cp_parser *parser, location_t clause_loc, tree list)
+{
+  tree c, expr = error_mark_node;
+  tree tile = NULL_TREE;
+
+  check_no_duplicate_clause (list, OMP_CLAUSE_TILE, "tile", clause_loc);
+
+  if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
+    return list;
+
+  do
+    {
+      if (cp_lexer_next_token_is (parser->lexer, CPP_MULT)
+	  && (cp_lexer_nth_token_is (parser->lexer, 2, CPP_COMMA)
+	      || cp_lexer_nth_token_is (parser->lexer, 2, CPP_CLOSE_PAREN)))
+	{
+	  cp_lexer_consume_token (parser->lexer);
+	  expr = integer_minus_one_node;
+	}
+      else
+	expr = cp_parser_assignment_expression (parser, NULL, false, false);
+
+      if (expr == error_mark_node)
+	return list;
+
+      tile = tree_cons (NULL_TREE, expr, tile);
+
+      if (cp_lexer_next_token_is (parser->lexer, CPP_COMMA))
+	cp_lexer_consume_token (parser->lexer);
+    }
+  while (cp_lexer_next_token_is_not (parser->lexer, CPP_CLOSE_PAREN));
+
+  /* Consume the trailing ')'.  */
+  cp_lexer_consume_token (parser->lexer);
+
+  c = build_omp_clause (clause_loc, OMP_CLAUSE_TILE);
+  tile = nreverse (tile);
+  OMP_CLAUSE_TILE_LIST (c) = tile;
+  OMP_CLAUSE_CHAIN (c) = list;
+  return c;
+}
+
 /* OpenACC 2.0
    Parse wait clause or directive parameters.  */
 
@@ -29859,10 +29912,14 @@ cp_parser_omp_clause_collapse (cp_parser *parser, tree list, location_t location
 }
 
 /* OpenMP 2.5:
-   default ( shared | none ) */
+   default ( shared | none )
+
+   OpenACC 2.0
+   default (none) */
 
 static tree
-cp_parser_omp_clause_default (cp_parser *parser, tree list, location_t location)
+cp_parser_omp_clause_default (cp_parser *parser, tree list,
+			      location_t location, bool is_oacc)
 {
   enum omp_clause_default_kind kind = OMP_CLAUSE_DEFAULT_UNSPECIFIED;
   tree c;
@@ -29883,7 +29940,7 @@ cp_parser_omp_clause_default (cp_parser *parser, tree list, location_t location)
 	  break;
 
 	case 's':
-	  if (strcmp ("shared", p) != 0)
+	  if (strcmp ("shared", p) != 0 || is_oacc)
 	    goto invalid_kind;
 	  kind = OMP_CLAUSE_DEFAULT_SHARED;
 	  break;
@@ -29897,7 +29954,10 @@ cp_parser_omp_clause_default (cp_parser *parser, tree list, location_t location)
   else
     {
     invalid_kind:
-      cp_parser_error (parser, "expected %<none%> or %<shared%>");
+      if (is_oacc)
+	cp_parser_error (parser, "expected %<none%>");
+      else
+	cp_parser_error (parser, "expected %<none%> or %<shared%>");
     }
 
   if (!cp_parser_require (parser, CPP_CLOSE_PAREN, RT_CLOSE_PAREN))
@@ -31444,6 +31504,10 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
 	  clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
 	  c_name = "delete";
 	  break;
+	case PRAGMA_OMP_CLAUSE_DEFAULT:
+	  clauses = cp_parser_omp_clause_default (parser, clauses, here, true);
+	  c_name = "default";
+	  break;
 	case PRAGMA_OACC_CLAUSE_DEVICE:
 	  clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
 	  c_name = "device";
@@ -31452,6 +31516,11 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
 	  clauses = cp_parser_oacc_data_clause_deviceptr (parser, clauses);
 	  c_name = "deviceptr";
 	  break;
+	case PRAGMA_OACC_CLAUSE_FIRSTPRIVATE:
+	  clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_FIRSTPRIVATE,
+					    clauses);
+	  c_name = "firstprivate";
+	  break;
 	case PRAGMA_OACC_CLAUSE_GANG:
 	  c_name = "gang";
 	  clauses = cp_parser_oacc_shape_clause (parser, OMP_CLAUSE_GANG,
@@ -31465,6 +31534,12 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
 	  clauses = cp_parser_omp_clause_if (parser, clauses, here, false);
 	  c_name = "if";
 	  break;
+	case PRAGMA_OACC_CLAUSE_INDEPENDENT:
+	  clauses = cp_parser_oacc_simple_clause (parser,
+						  OMP_CLAUSE_INDEPENDENT,
+						  clauses, here);
+	  c_name = "independent";
+	  break;
 	case PRAGMA_OACC_CLAUSE_NUM_GANGS:
 	  code = OMP_CLAUSE_NUM_GANGS;
 	  c_name = "num_gangs";
@@ -31497,6 +31572,11 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
 	  clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
 	  c_name = "present_or_create";
 	  break;
+	case PRAGMA_OACC_CLAUSE_PRIVATE:
+	  clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_PRIVATE,
+					    clauses);
+	  c_name = "private";
+	  break;
 	case PRAGMA_OACC_CLAUSE_REDUCTION:
 	  clauses = cp_parser_omp_clause_reduction (parser, clauses);
 	  c_name = "reduction";
@@ -31510,6 +31590,10 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
 						 clauses, here);
 	  c_name = "seq";
 	  break;
+	case PRAGMA_OACC_CLAUSE_TILE:
+	  clauses = cp_parser_oacc_clause_tile (parser, here, clauses);
+	  c_name = "tile";
+	  break;
 	case PRAGMA_OACC_CLAUSE_VECTOR:
 	  c_name = "vector";
 	  clauses = cp_parser_oacc_shape_clause (parser, OMP_CLAUSE_VECTOR,
@@ -31598,7 +31682,7 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask,
 	  break;
 	case PRAGMA_OMP_CLAUSE_DEFAULT:
 	  clauses = cp_parser_omp_clause_default (parser, clauses,
-						  token->location);
+						  token->location, false);
 	  c_name = "default";
 	  break;
 	case PRAGMA_OMP_CLAUSE_FINAL:
@@ -34493,12 +34577,15 @@ cp_parser_oacc_enter_exit_data (cp_parser *parser, cp_token *pragma_tok,
 
 #define OACC_LOOP_CLAUSE_MASK						\
 	( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COLLAPSE)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRIVATE)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_GANG)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WORKER)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_AUTO)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_INDEPENDENT)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ)			\
-	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) )
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_TILE))
 
 static tree
 cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
@@ -34543,6 +34630,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
 	| (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_DEVICEPTR)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF)			\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT)		\
@@ -34558,7 +34646,9 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
 	| (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_DEVICEPTR)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FIRSTPRIVATE)       	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF)			\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS)		\
@@ -34567,6 +34657,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE)   \
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRIVATE)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH)       \
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
@@ -34642,7 +34733,7 @@ cp_parser_oacc_update (cp_parser *parser, cp_token *pragma_tok)
     {
       error_at (pragma_tok->location,
 		"%<#pragma acc update%> must contain at least one "
-		"%<device%> or %<host/self%> clause");
+		"%<device%> or %<host%> or %<self%> clause");
       return NULL_TREE;
     }
 
diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c
index 9e3bd2d..45eda3a 100644
--- a/gcc/cp/pt.c
+++ b/gcc/cp/pt.c
@@ -14395,6 +14395,13 @@ tsubst_omp_clauses (tree clauses, bool declare_simd, bool allow_fields,
 	case OMP_CLAUSE_PRIORITY:
 	case OMP_CLAUSE_ORDERED:
 	case OMP_CLAUSE_HINT:
+	case OMP_CLAUSE_NUM_GANGS:
+	case OMP_CLAUSE_NUM_WORKERS:
+	case OMP_CLAUSE_VECTOR_LENGTH:
+	case OMP_CLAUSE_WORKER:
+	case OMP_CLAUSE_VECTOR:
+	case OMP_CLAUSE_ASYNC:
+	case OMP_CLAUSE_WAIT:
 	  OMP_CLAUSE_OPERAND (nc, 0)
 	    = tsubst_expr (OMP_CLAUSE_OPERAND (oc, 0), args, complain, 
 			   in_decl, /*integral_constant_expression_p=*/false);
@@ -14419,6 +14426,7 @@ tsubst_omp_clauses (tree clauses, bool declare_simd, bool allow_fields,
 	    = tsubst_omp_clause_decl (OMP_CLAUSE_DECL (oc), args, complain,
 				      in_decl);
 	  break;
+	case OMP_CLAUSE_GANG:
 	case OMP_CLAUSE_ALIGNED:
 	  OMP_CLAUSE_DECL (nc)
 	    = tsubst_omp_clause_decl (OMP_CLAUSE_DECL (oc), args, complain,
@@ -14461,6 +14469,22 @@ tsubst_omp_clauses (tree clauses, bool declare_simd, bool allow_fields,
 	case OMP_CLAUSE_THREADS:
 	case OMP_CLAUSE_SIMD:
 	case OMP_CLAUSE_DEFAULTMAP:
+	case OMP_CLAUSE_INDEPENDENT:
+	case OMP_CLAUSE_AUTO:
+	case OMP_CLAUSE_SEQ:
+	  break;
+	case OMP_CLAUSE_TILE:
+	  {
+	    tree lnc, loc;
+	    for (lnc = OMP_CLAUSE_TILE_LIST (nc),
+		   loc = OMP_CLAUSE_TILE_LIST (oc);
+		 loc;
+		 loc = TREE_CHAIN (loc), lnc = TREE_CHAIN (lnc))
+	      {
+		TREE_VALUE (lnc) = tsubst_expr (TREE_VALUE (loc), args,
+						complain, in_decl, false);
+	      }
+	  }
 	  break;
 	default:
 	  gcc_unreachable ();
@@ -15231,6 +15255,15 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
       }
       break;
 
+    case OACC_KERNELS:
+    case OACC_PARALLEL:
+      tmp = tsubst_omp_clauses (OMP_CLAUSES (t), false, false, args, complain,
+				in_decl);
+      stmt = begin_omp_parallel ();
+      RECUR (OMP_BODY (t));
+      finish_omp_construct (TREE_CODE (t), stmt, tmp);
+      break;
+
     case OMP_PARALLEL:
       r = push_omp_privatization_clauses (OMP_PARALLEL_COMBINED (t));
       tmp = tsubst_omp_clauses (OMP_PARALLEL_CLAUSES (t), false, true,
@@ -15261,6 +15294,7 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
     case CILK_FOR:
     case OMP_DISTRIBUTE:
     case OMP_TASKLOOP:
+    case OACC_LOOP:
       {
 	tree clauses, body, pre_body;
 	tree declv = NULL_TREE, initv = NULL_TREE, condv = NULL_TREE;
@@ -15269,7 +15303,8 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
 	int i;
 
 	r = push_omp_privatization_clauses (OMP_FOR_INIT (t) == NULL_TREE);
-	clauses = tsubst_omp_clauses (OMP_FOR_CLAUSES (t), false, true,
+	clauses = tsubst_omp_clauses (OMP_FOR_CLAUSES (t), false,
+				      TREE_CODE (t) != OACC_LOOP,
 				      args, complain, in_decl);
 	if (OMP_FOR_INIT (t) != NULL_TREE)
 	  {
@@ -15339,9 +15374,11 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
       pop_omp_privatization_clauses (r);
       break;
 
+    case OACC_DATA:
     case OMP_TARGET_DATA:
     case OMP_TARGET:
-      tmp = tsubst_omp_clauses (OMP_CLAUSES (t), false, true,
+      tmp = tsubst_omp_clauses (OMP_CLAUSES (t), false,
+				TREE_CODE (t) != OACC_DATA,
 				args, complain, in_decl);
       keep_next_level (true);
       stmt = begin_omp_structured_block ();
@@ -15395,6 +15432,16 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
       add_stmt (t);
       break;
 
+    case OACC_ENTER_DATA:
+    case OACC_EXIT_DATA:
+    case OACC_UPDATE:
+      tmp = tsubst_omp_clauses (OMP_STANDALONE_CLAUSES (t), false, false,
+				args, complain, in_decl);
+      t = copy_node (t);
+      OMP_STANDALONE_CLAUSES (t) = tmp;
+      add_stmt (t);
+      break;
+
     case OMP_ORDERED:
       tmp = tsubst_omp_clauses (OMP_ORDERED_CLAUSES (t), false, true,
 				args, complain, in_decl);
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index fe18b27..7702a41 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -6855,9 +6855,47 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd)
 	case OMP_CLAUSE_DEFAULTMAP:
 	case OMP_CLAUSE__CILK_FOR_COUNT_:
 	case OMP_CLAUSE_AUTO:
+	case OMP_CLAUSE_INDEPENDENT:
 	case OMP_CLAUSE_SEQ:
 	  break;
 
+	case OMP_CLAUSE_TILE:
+	  for (tree list = OMP_CLAUSE_TILE_LIST (c); !remove && list;
+	       list = TREE_CHAIN (list))
+	    {
+	      t = TREE_VALUE (list);
+
+	      if (t == error_mark_node)
+		remove = true;
+	      else if (!type_dependent_expression_p (t)
+		       && !INTEGRAL_TYPE_P (TREE_TYPE (t)))
+		{
+		  error ("%<tile%> value must be integral");
+		  remove = true;
+		}
+	      else
+		{
+		  t = mark_rvalue_use (t);
+		  if (!processing_template_decl)
+		    {
+		      t = maybe_constant_value (t);
+		      if (TREE_CODE (t) == INTEGER_CST
+			  && tree_int_cst_sgn (t) != 1
+			  && t != integer_minus_one_node)
+			{
+			  warning_at (OMP_CLAUSE_LOCATION (c), 0,
+				      "%<tile%> value must be positive");
+			  t = integer_one_node;
+			}
+		    }
+		  t = fold_build_cleanup_point_expr (TREE_TYPE (t), t);
+		}
+
+		/* Update list item.  */
+	      TREE_VALUE (list) = t;
+	    }
+	  break;
+
 	case OMP_CLAUSE_ORDERED:
 	  ordered_seen = true;
 	  break;
diff --git a/gcc/fortran/ChangeLog b/gcc/fortran/ChangeLog
index e2c8f71..bd3fa76 100644
--- a/gcc/fortran/ChangeLog
+++ b/gcc/fortran/ChangeLog
@@ -1,3 +1,17 @@
+2015-11-05  Cesar Philippidis  <cesar@codesourcery.com>
+
+	* openmp.c (gfc_match_omp_clauses): Update support for the tile
+	and default clauses in OpenACC.
+	(gfc_match_oacc_update): Error when data clauses are supplied.
+	(oacc_compatible_clauses): Delete.
+	(resolve_omp_clauses): Give special care for OpenACC reductions.
+	Also update error reporting for the tile clause.
+	(resolve_oacc_loop_blocks): Update error reporting for the tile clause.
+	* trans-openmp.c (gfc_trans_omp_clauses): Update OMP_CLAUSE_SEQ. Add
+	OMP_CLAUSE_{AUTO,TILE} and add support the the gang static argument.
+	(gfc_trans_oacc_combined_directive): Update the list of clauses which
+	are split to acc loops.
+
 2015-11-05  Jakub Jelinek  <jakub@redhat.com>
 
 	* types.def (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR): Remove.
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index 929a739..a7c7a19 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -703,6 +703,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
 				       OMP_MAP_FORCE_FROM))
 	continue;
       if ((mask & OMP_CLAUSE_TILE)
+	  && !c->tile_list
 	  && match_oacc_expr_list ("tile (", &c->tile_list, true) == MATCH_YES)
 	continue;
       if ((mask & OMP_CLAUSE_SEQ) && !c->seq
@@ -856,12 +857,14 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
       if ((mask & OMP_CLAUSE_DEFAULT)
 	  && c->default_sharing == OMP_DEFAULT_UNKNOWN)
 	{
-	  if (gfc_match ("default ( shared )") == MATCH_YES)
+	  if (gfc_match ("default ( none )") == MATCH_YES)
+	    c->default_sharing = OMP_DEFAULT_NONE;
+	  else if (openacc)
+	    /* c->default_sharing = OMP_DEFAULT_UNKNOWN */;
+	  else if (gfc_match ("default ( shared )") == MATCH_YES)
 	    c->default_sharing = OMP_DEFAULT_SHARED;
 	  else if (gfc_match ("default ( private )") == MATCH_YES)
 	    c->default_sharing = OMP_DEFAULT_PRIVATE;
-	  else if (gfc_match ("default ( none )") == MATCH_YES)
-	    c->default_sharing = OMP_DEFAULT_NONE;
 	  else if (gfc_match ("default ( firstprivate )") == MATCH_YES)
 	    c->default_sharing = OMP_DEFAULT_FIRSTPRIVATE;
 	  if (c->default_sharing != OMP_DEFAULT_UNKNOWN)
@@ -1304,10 +1307,19 @@ match
 gfc_match_oacc_update (void)
 {
   gfc_omp_clauses *c;
+  locus here = gfc_current_locus;
+
   if (gfc_match_omp_clauses (&c, OACC_UPDATE_CLAUSES, false, false, true)
       != MATCH_YES)
     return MATCH_ERROR;
 
+  if (!c->lists[OMP_LIST_MAP])
+    {
+      gfc_error ("%<acc update%> must contain at least one "
+		 "%<device%> or %<host%> or %<self%> clause at %L", &here);
+      return MATCH_ERROR;
+    }
+
   new_st.op = EXEC_OACC_UPDATE;
   new_st.ext.omp_clauses = c;
   return MATCH_YES;
@@ -2846,30 +2858,6 @@ resolve_omp_udr_clause (gfc_omp_namelist *n, gfc_namespace *ns,
   return copy;
 }
 
-/* Returns true if clause in list 'list' is compatible with any of
-   of the clauses in lists [0..list-1].  E.g., a reduction variable may
-   appear in both reduction and private clauses, so this function
-   will return true in this case.  */
-
-static bool
-oacc_compatible_clauses (gfc_omp_clauses *clauses, int list,
-			   gfc_symbol *sym, bool openacc)
-{
-  gfc_omp_namelist *n;
-
-  if (!openacc)
-    return false;
-
-  if (list != OMP_LIST_REDUCTION)
-    return false;
-
-  for (n = clauses->lists[OMP_LIST_FIRST]; n; n = n->next)
-    if (n->sym == sym)
-      return true;
-
-  return false;
-}
-
 /* OpenMP directive resolving routines.  */
 
 static void
@@ -2975,11 +2963,11 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 	&& list != OMP_LIST_DEPEND
 	&& (list != OMP_LIST_MAP || openacc)
 	&& list != OMP_LIST_FROM
-	&& list != OMP_LIST_TO)
+	&& list != OMP_LIST_TO
+	&& (list != OMP_LIST_REDUCTION || !openacc))
       for (n = omp_clauses->lists[list]; n; n = n->next)
 	{
-	  if (n->sym->mark && !oacc_compatible_clauses (omp_clauses, list,
-							n->sym, openacc))
+	  if (n->sym->mark)
 	    gfc_error ("Symbol %qs present on multiple clauses at %L",
 		       n->sym->name, &n->where);
 	  else
@@ -3028,6 +3016,22 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 	n->sym->mark = 1;
     }
 
+  /* OpenACC reductions.  */
+  if (openacc)
+    {
+      for (n = omp_clauses->lists[OMP_LIST_REDUCTION]; n; n = n->next)
+	n->sym->mark = 0;
+
+      for (n = omp_clauses->lists[OMP_LIST_REDUCTION]; n; n = n->next)
+	{
+	  if (n->sym->mark)
+	    gfc_error ("Symbol %qs present on multiple clauses at %L",
+		       n->sym->name, &n->where);
+	  else
+	    n->sym->mark = 1;
+	}
+    }
+  
   for (n = omp_clauses->lists[OMP_LIST_TO]; n; n = n->next)
     n->sym->mark = 0;
   for (n = omp_clauses->lists[OMP_LIST_FROM]; n; n = n->next)
@@ -4528,22 +4532,8 @@ resolve_oacc_loop_blocks (gfc_code *code)
       if (code->ext.omp_clauses->vector)
 	gfc_error ("Clause AUTO conflicts with VECTOR at %L", &code->loc);
     }
-  if (!code->ext.omp_clauses->tile_list)
-    {
-      if (code->ext.omp_clauses->gang)
-	{
-	  if (code->ext.omp_clauses->worker)
-	    gfc_error ("Clause GANG conflicts with WORKER at %L", &code->loc);
-	  if (code->ext.omp_clauses->vector)
-	    gfc_error ("Clause GANG conflicts with VECTOR at %L", &code->loc);
-	}
-      if (code->ext.omp_clauses->worker)
-	if (code->ext.omp_clauses->vector)
-	  gfc_error ("Clause WORKER conflicts with VECTOR at %L", &code->loc);
-    }
-  else if (code->ext.omp_clauses->gang
-	   && code->ext.omp_clauses->worker
-	   && code->ext.omp_clauses->vector)
+  if (code->ext.omp_clauses->tile_list && code->ext.omp_clauses->gang
+      && code->ext.omp_clauses->worker && code->ext.omp_clauses->vector)
     gfc_error ("Tiled loop cannot be parallelized across gangs, workers and "
 	       "vectors at the same time at %L", &code->loc);
 
@@ -4564,10 +4554,21 @@ resolve_oacc_loop_blocks (gfc_code *code)
 	{
 	  num++;
 	  if (el->expr == NULL)
-	    continue;
-	  resolve_oacc_positive_int_expr (el->expr, "TILE");
-	  if (el->expr->expr_type != EXPR_CONSTANT)
-	    gfc_error ("TILE requires constant expression at %L", &code->loc);
+	    {
+	      /* NULL expressions are used to represent '*' arguments.
+		 Convert those to a -1 expressions.  */
+	      el->expr = gfc_get_constant_expr (BT_INTEGER,
+						gfc_default_integer_kind,
+						&code->loc);
+	      mpz_set_si (el->expr->value.integer, -1);
+	    }
+	  else
+	    {
+	      resolve_oacc_positive_int_expr (el->expr, "TILE");
+	      if (el->expr->expr_type != EXPR_CONSTANT)
+		gfc_error ("TILE requires constant expression at %L",
+			   &code->loc);
+	    }
 	}
       resolve_oacc_nested_loops (code, code->block->next, num, "tiled");
     }
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index 5f4c382..9f0d533 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -2534,8 +2534,12 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
     }
   if (clauses->seq)
     {
-      c = build_omp_clause (where.lb->location, OMP_CLAUSE_ORDERED);
-      OMP_CLAUSE_ORDERED_EXPR (c) = NULL_TREE;
+      c = build_omp_clause (where.lb->location, OMP_CLAUSE_SEQ);
+      omp_clauses = gfc_trans_add_clause (c, omp_clauses);
+    }
+  if (clauses->par_auto)
+    {
+      c = build_omp_clause (where.lb->location, OMP_CLAUSE_AUTO);
       omp_clauses = gfc_trans_add_clause (c, omp_clauses);
     }
   if (clauses->independent)
@@ -2579,6 +2583,21 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
       OMP_CLAUSE_VECTOR_LENGTH_EXPR (c) = vector_length_var;
       omp_clauses = gfc_trans_add_clause (c, omp_clauses);
     }
+  if (clauses->tile_list)
+    {
+      vec<tree, va_gc> *tvec;
+      gfc_expr_list *el;
+
+      vec_alloc (tvec, 4);
+
+      for (el = clauses->tile_list; el; el = el->next)
+	vec_safe_push (tvec, gfc_convert_expr_to_tree (block, el->expr));
+
+      c = build_omp_clause (where.lb->location, OMP_CLAUSE_TILE);
+      OMP_CLAUSE_TILE_LIST (c) = build_tree_list_vec (tvec);
+      omp_clauses = gfc_trans_add_clause (c, omp_clauses);
+      tvec->truncate (0);
+    }
   if (clauses->vector)
     {
       if (clauses->vector_expr)
@@ -2618,7 +2637,17 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 	  tree gang_var
 	    = gfc_convert_expr_to_tree (block, clauses->gang_expr);
 	  c = build_omp_clause (where.lb->location, OMP_CLAUSE_GANG);
-	  OMP_CLAUSE_GANG_EXPR (c) = gang_var;
+	  if (clauses->gang_static)
+	    OMP_CLAUSE_GANG_STATIC_EXPR (c) = gang_var;
+	  else
+	    OMP_CLAUSE_GANG_EXPR (c) = gang_var;
+	  omp_clauses = gfc_trans_add_clause (c, omp_clauses);
+	}
+      else if (clauses->gang_static)
+	{
+	  /* This corresponds to gang (static: *).  */
+	  c = build_omp_clause (where.lb->location, OMP_CLAUSE_GANG);
+	  OMP_CLAUSE_GANG_STATIC_EXPR (c) = integer_minus_one_node;
 	  omp_clauses = gfc_trans_add_clause (c, omp_clauses);
 	}
       else
@@ -3449,16 +3478,33 @@ gfc_trans_oacc_combined_directive (gfc_code *code)
 	      sizeof (construct_clauses));
       loop_clauses.collapse = construct_clauses.collapse;
       loop_clauses.gang = construct_clauses.gang;
+      loop_clauses.gang_expr = construct_clauses.gang_expr;
+      loop_clauses.gang_static = construct_clauses.gang_static;
       loop_clauses.vector = construct_clauses.vector;
+      loop_clauses.vector_expr = construct_clauses.vector_expr;
       loop_clauses.worker = construct_clauses.worker;
+      loop_clauses.worker_expr = construct_clauses.worker_expr;
       loop_clauses.seq = construct_clauses.seq;
+      loop_clauses.par_auto = construct_clauses.par_auto;
       loop_clauses.independent = construct_clauses.independent;
-      construct_clauses.collapse = 0;
+      loop_clauses.tile_list = construct_clauses.tile_list;
+      loop_clauses.lists[OMP_LIST_PRIVATE]
+	= construct_clauses.lists[OMP_LIST_PRIVATE];
+      loop_clauses.lists[OMP_LIST_REDUCTION]
+	= construct_clauses.lists[OMP_LIST_REDUCTION];
       construct_clauses.gang = false;
+      construct_clauses.gang_expr = NULL;
+      construct_clauses.gang_static = false;
       construct_clauses.vector = false;
+      construct_clauses.vector_expr = NULL;
       construct_clauses.worker = false;
+      construct_clauses.worker_expr = NULL;
       construct_clauses.seq = false;
+      construct_clauses.par_auto = false;
+      construct_clauses.independent = false;
       construct_clauses.independent = false;
+      construct_clauses.tile_list = NULL;
+      construct_clauses.lists[OMP_LIST_PRIVATE] = NULL;
       oacc_clauses = gfc_trans_omp_clauses (&block, &construct_clauses,
 					    code->loc);
     }
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 319e200..fa34858 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -7141,9 +7141,18 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	    remove = true;
 	  break;
 
+	case OMP_CLAUSE_TILE:
+	  for (tree list = OMP_CLAUSE_TILE_LIST (c); !remove && list;
+	       list = TREE_CHAIN (list))
+	    {
+	      if (gimplify_expr (&TREE_VALUE (list), pre_p, NULL,
+				 is_gimple_val, fb_rvalue) == GS_ERROR)
+		remove = true;
+	    }
+	  break;
+
 	case OMP_CLAUSE_DEVICE_RESIDENT:
 	case OMP_CLAUSE_USE_DEVICE:
-	case OMP_CLAUSE_INDEPENDENT:
 	  remove = true;
 	  break;
 
@@ -7153,6 +7162,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	case OMP_CLAUSE_COLLAPSE:
 	case OMP_CLAUSE_AUTO:
 	case OMP_CLAUSE_SEQ:
+	case OMP_CLAUSE_INDEPENDENT:
 	case OMP_CLAUSE_MERGEABLE:
 	case OMP_CLAUSE_PROC_BIND:
 	case OMP_CLAUSE_SAFELEN:
@@ -7663,6 +7673,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, tree *list_p,
 	case OMP_CLAUSE_VECTOR:
 	case OMP_CLAUSE_AUTO:
 	case OMP_CLAUSE_SEQ:
+	case OMP_CLAUSE_TILE:
 	  break;
 
 	default:
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 86cc655..5ffb276 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -2129,6 +2129,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	case OMP_CLAUSE_GANG:
 	case OMP_CLAUSE_WORKER:
 	case OMP_CLAUSE_VECTOR:
+	case OMP_CLAUSE_TILE:
 	  break;
 
 	case OMP_CLAUSE_ALIGNED:
@@ -2309,6 +2310,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	case OMP_CLAUSE_GANG:
 	case OMP_CLAUSE_WORKER:
 	case OMP_CLAUSE_VECTOR:
+	case OMP_CLAUSE_TILE:
 	  break;
 
 	case OMP_CLAUSE_DEVICE_RESIDENT:
diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog
index eb51808..c074e7b 100644
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,3 +1,32 @@
+2015-11-05  Cesar Philippidis  <cesar@codesourcery.com>
+	    Tom de Vries  <tom@codesourcery.com>
+	    Nathan Sidwell  <nathan@codesourcery.com>
+	    Thomas Schwinge  <thomas@codesourcery.com>
+
+	* c-c++-common/goacc/combined-directives.c: New test.
+	* c-c++-common/goacc/loop-clauses.c: New test.
+	* c-c++-common/goacc/tile.c: New test.
+	* c-c++-common/goacc/loop-shape.c: Add test for pointer variable
+	as gang static arguments.
+	* c-c++-common/goacc/update-1.c: Adjust expected error message.
+	* g++.dg/goacc/template.C: New test.
+	* gfortran.dg/goacc/combined-directives.f90: New test.
+	* gfortran.dg/goacc/default.f95: New test.
+	* gfortran.dg/goacc/default_none.f95: New test.
+	* gfortran.dg/goacc/firstprivate-1.f95: New test.
+	* gfortran.dg/goacc/gang-static.f95: New test.
+	* gfortran.dg/goacc/kernels-loop-inner.f95: New test.
+	* gfortran.dg/goacc/kernels-loops-adjacent.f95: New test.
+	* gfortran.dg/goacc/list.f95: Update test.
+	* gfortran.dg/goacc/loop-2.f95: Likewise.
+	* gfortran.dg/goacc/loop-4.f95: New test.
+	* gfortran.dg/goacc/loop-5.f95: New test.
+	* gfortran.dg/goacc/loop-6.f95: New test.
+	* gfortran.dg/goacc/loop-tree-1.f90: Update test.
+	* gfortran.dg/goacc/multi-clause.f90: New test.
+	* gfortran.dg/goacc/parallel-tree.f95: Update test.
+	* gfortran.dg/goacc/update.f95: New test.
+
 2015-11-05  Martin Sebor  <msebor@redhat.com>
 
 	PR c++/67942
diff --git a/gcc/testsuite/c-c++-common/goacc/combined-directives.c b/gcc/testsuite/c-c++-common/goacc/combined-directives.c
new file mode 100644
index 0000000..c387285
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/combined-directives.c
@@ -0,0 +1,119 @@
+// { dg-do compile }
+// { dg-options "-fopenacc -fdump-tree-gimple" }
+
+// This error is temporary.  Remove when support is added for these clauses
+// in the middle end.  Also remove the comments from the reduction test
+// after the FE learns that reduction variables may appear in data clauses too.
+// { dg-prune-output "sorry, unimplemented" }
+
+void
+test ()
+{
+  int a[100], i, j, z;
+
+  // acc parallel
+
+  #pragma acc parallel loop collapse (2)
+  for (i = 0; i < 100; i++)
+    for (j = 0; j < 10; j++)
+      ;
+
+  #pragma acc parallel loop gang
+  for (i = 0; i < 100; i++)
+    ;
+
+  #pragma acc parallel loop worker
+  for (i = 0; i < 100; i++)
+    for (j = 0; j < 10; j++)
+      ;
+
+  #pragma acc parallel loop vector
+  for (i = 0; i < 100; i++)
+    for (j = 0; j < 10; j++)
+      ;
+
+  #pragma acc parallel loop seq
+  for (i = 0; i < 100; i++)
+    for (j = 0; j < 10; j++)
+      ;
+
+  #pragma acc parallel loop auto
+  for (i = 0; i < 100; i++)
+    for (j = 0; j < 10; j++)
+      ;
+
+  #pragma acc parallel loop tile (2, 3)
+  for (i = 0; i < 100; i++)
+    for (j = 0; j < 10; j++)
+      ;
+
+  #pragma acc parallel loop independent
+  for (i = 0; i < 100; i++)
+    ;
+
+  #pragma acc parallel loop private (z)
+  for (i = 0; i < 100; i++)
+    z = 0;
+
+//  #pragma acc parallel loop reduction (+:z) copy (z)
+//  for (i = 0; i < 100; i++)
+//    ;
+
+  // acc kernels
+
+  #pragma acc kernels loop collapse (2)
+  for (i = 0; i < 100; i++)
+    for (j = 0; j < 10; j++)
+      ;
+
+  #pragma acc kernels loop gang
+  for (i = 0; i < 100; i++)
+    ;
+
+  #pragma acc kernels loop worker
+  for (i = 0; i < 100; i++)
+    for (j = 0; j < 10; j++)
+      ;
+
+  #pragma acc kernels loop vector
+  for (i = 0; i < 100; i++)
+    for (j = 0; j < 10; j++)
+      ;
+
+  #pragma acc kernels loop seq
+  for (i = 0; i < 100; i++)
+    for (j = 0; j < 10; j++)
+      ;
+
+  #pragma acc kernels loop auto
+  for (i = 0; i < 100; i++)
+    for (j = 0; j < 10; j++)
+      ;
+
+  #pragma acc kernels loop tile (2, 3)
+  for (i = 0; i < 100; i++)
+    for (j = 0; j < 10; j++)
+      ;
+
+  #pragma acc kernels loop independent
+  for (i = 0; i < 100; i++)
+    ;
+
+  #pragma acc kernels loop private (z)
+  for (i = 0; i < 100; i++)
+    z = 0;
+
+//  #pragma acc kernels loop reduction (+:z) copy (z)
+//  for (i = 0; i < 100; i++)
+//    ;
+}
+
+// { dg-final { scan-tree-dump-times "acc loop collapse.2. private.j. private.i" 2 "gimple" } }
+// { dg-final { scan-tree-dump-times "acc loop gang" 2 "gimple" } }
+// { dg-final { scan-tree-dump-times "acc loop worker" 2 "gimple" } }
+// { dg-final { scan-tree-dump-times "acc loop vector" 2 "gimple" } }
+// { dg-final { scan-tree-dump-times "acc loop seq" 2 "gimple" } }
+// { dg-final { scan-tree-dump-times "acc loop auto" 2 "gimple" } }
+// { dg-final { scan-tree-dump-times "acc loop tile.2, 3" 2 "gimple" } }
+// { dg-final { scan-tree-dump-times "acc loop independent private.i" 2 "gimple" } }
+// { dg-final { scan-tree-dump-times "private.z" 2 "gimple" } }
diff --git a/gcc/testsuite/c-c++-common/goacc/loop-clauses.c b/gcc/testsuite/c-c++-common/goacc/loop-clauses.c
new file mode 100644
index 0000000..97b8786
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/loop-clauses.c
@@ -0,0 +1,86 @@
+/* { dg-do compile } */
+
+/* { dg-prune-output "sorry, unimplemented" } */
+
+int
+main ()
+{
+  int i, j;
+
+#pragma acc parallel firstprivate (j) private (i)
+  {
+#pragma acc loop seq
+    for (i = 0; i < 10; i++)
+      { }
+  }
+
+#pragma acc parallel default (none)
+  {
+#pragma acc loop auto private (j)
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop gang
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop gang(static:5)
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop gang(static:*)
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop vector
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop worker
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop auto
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop independent
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop seq
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop gang worker vector
+    for (i = 0; i < 10; i++)
+      { }
+  }
+
+#pragma acc kernels default (none)
+  {
+#pragma acc loop auto
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop gang (num:5)
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop gang(static:5)
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop gang(static:*)
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop vector(length:10)
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop worker(num:5)
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop auto
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop independent
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop seq
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop gang worker vector
+    for (i = 0; i < 10; i++)
+      { }
+  }
+
+  return 0;
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/loop-shape.c b/gcc/testsuite/c-c++-common/goacc/loop-shape.c
index b6d3156..9708f7b 100644
--- a/gcc/testsuite/c-c++-common/goacc/loop-shape.c
+++ b/gcc/testsuite/c-c++-common/goacc/loop-shape.c
@@ -8,6 +8,7 @@ int main ()
   int i;
   int v = 32, w = 19;
   int length = 1, num = 5;
+  int *abc;
 
   /* Valid uses.  */
 
@@ -199,12 +200,12 @@ int main ()
     ;
 
   #pragma acc kernels
-  #pragma acc loop gang(static: * abc) /* { dg-error "expected '.' before" } */
+  #pragma acc loop gang(static: * abc)
   for (i = 0; i < 10; i++)
     ;
 
   #pragma acc kernels
-  #pragma acc loop gang(static:*num:1) /* { dg-error "expected '.' before" } */
+  #pragma acc loop gang(static:*num:1) /* { dg-error "" } */
   for (i = 0; i < 10; i++)
     ;
 
diff --git a/gcc/testsuite/c-c++-common/goacc/tile.c b/gcc/testsuite/c-c++-common/goacc/tile.c
new file mode 100644
index 0000000..2a81427
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/tile.c
@@ -0,0 +1,73 @@
+/* { dg-do compile } */
+
+int
+main ()
+{
+  int i, *a, b;
+
+#pragma acc parallel loop tile (10)
+  for (i = 0; i < 100; i++)
+    ;
+
+#pragma acc parallel loop tile (*)
+  for (i = 0; i < 100; i++)
+    ;
+
+#pragma acc parallel loop tile (10, *)
+  for (i = 0; i < 100; i++)
+    ;
+
+#pragma acc parallel loop tile (10, *, i)
+  for (i = 0; i < 100; i++)
+    ;
+
+#pragma acc parallel loop tile // { dg-error "expected '\\\('" }
+  for (i = 0; i < 100; i++)
+    ;  
+
+#pragma acc parallel loop tile () // { dg-error "" }
+  for (i = 0; i < 100; i++)
+    ;
+
+#pragma acc parallel loop tile (,1) // { dg-error "" }
+  for (i = 0; i < 100; i++)
+    ;
+
+#pragma acc parallel loop tile (,,) // { dg-error "" }
+  for (i = 0; i < 100; i++)
+    ;
+
+#pragma acc parallel loop tile (1.1) // { dg-error "'tile' value must be integral" }
+  for (i = 0; i < 100; i++)
+    ;
+
+#pragma acc parallel loop tile (-3) // { dg-warning "'tile' value must be positive" }
+  for (i = 0; i < 100; i++)
+    ;
+
+#pragma acc parallel loop tile (10,-3) // { dg-warning "'tile' value must be positive" }
+  for (i = 0; i < 100; i++)
+    ;
+
+#pragma acc parallel loop tile (-100,10,5) // { dg-warning "'tile' value must be positive" }
+  for (i = 0; i < 100; i++)
+    ;
+
+#pragma acc parallel loop tile (1,2.0,true) // { dg-error "" }
+  for (i = 0; i < 100; i++)
+    ;
+
+#pragma acc parallel loop tile (*a, 1)
+  for (i = 0; i < 100; i++)
+    ;
+
+#pragma acc parallel loop tile (1, *a, b)
+  for (i = 0; i < 100; i++)
+    ;
+
+#pragma acc parallel loop tile (b, 1, *a)
+  for (i = 0; i < 100; i++)
+    ;
+
+  return 0;
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/update-1.c b/gcc/testsuite/c-c++-common/goacc/update-1.c
index 97e9379..701ef36 100644
--- a/gcc/testsuite/c-c++-common/goacc/update-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/update-1.c
@@ -1,7 +1,7 @@
 void
 f (void)
 {
-#pragma acc update /* { dg-error "'#pragma acc update' must contain at least one 'device' or 'host/self' clause" } */
+#pragma acc update /* { dg-error "'#pragma acc update' must contain at least one 'device' or 'host' or 'self' clause" } */
 
   int i = 0;
   int a[10];
diff --git a/gcc/testsuite/g++.dg/goacc/template.C b/gcc/testsuite/g++.dg/goacc/template.C
new file mode 100644
index 0000000..f7a717b
--- /dev/null
+++ b/gcc/testsuite/g++.dg/goacc/template.C
@@ -0,0 +1,141 @@
+// This error is temporary.  Remove when support is added for these clauses
+// in the middle end.  Also remove the comments from the reduction test
+// after the FE learns that reduction variables may appear in data clauses too.
+// { dg-prune-output "sorry, unimplemented" }
+
+#pragma acc routine
+template <typename T> T
+accDouble(int val)
+{
+  return val * 2;
+}
+
+template<typename T> T
+oacc_parallel_copy (T a)
+{
+  T b = 0;
+  char w = 1;
+  int x = 2;
+  float y = 3;
+  double z = 4;
+
+#pragma acc parallel num_gangs (a) num_workers (a) vector_length (a) default (none) copyout (b) copyin (a)
+  {
+    b = a;
+  }
+
+#pragma acc parallel num_gangs (a) copy (w, x, y, z)
+  {
+    w = accDouble<char>(w);
+    x = accDouble<int>(x);
+    y = accDouble<float>(y);
+    z = accDouble<double>(z);
+  }
+
+#pragma acc parallel num_gangs (a) if (1)
+  {
+#pragma acc loop auto tile (a, 3)
+  for (int i = 0; i < a; i++)
+    for (int j = 0; j < 5; j++)
+      b = a;
+
+#pragma acc loop seq
+  for (int i = 0; i < a; i++)
+    b = a;
+  }
+
+  T c;
+
+#pragma acc parallel num_workers (10)
+  {
+#pragma acc atomic capture
+    c = b++;
+
+#pragma atomic update
+    c++;
+
+#pragma acc atomic read
+    b = a;
+
+#pragma acc atomic write
+    b = a;
+  }
+
+//#pragma acc parallel reduction (+:c)
+//  {
+//    c = 1;
+//  }
+
+#pragma acc data if (1) copy (b)
+  {
+    #pragma acc parallel
+    {
+      b = a;
+    }
+  }
+
+#pragma acc enter data copyin (b)
+#pragma acc parallel present (b)
+    {
+      b = a;
+    }
+
+#pragma acc update host (b)
+#pragma acc update self (b)
+#pragma acc update device (b)
+#pragma acc exit data delete (b)
+
+  return b;
+}
+
+template<typename T> T
+oacc_kernels_copy (T a)
+{
+  T b = 0;
+  T c = 0;
+  char w = 1;
+  int x = 2;
+  float y = 3;
+  double z = 4;
+
+#pragma acc kernels copy (w, x, y, z)
+  {
+    w = accDouble<char>(w);
+    x = accDouble<int>(x);
+    y = accDouble<float>(y);
+    z = accDouble<double>(z);
+  }
+
+#pragma acc kernels copyout (b) copyin (a)
+  b = a;
+
+//#pragma acc kernels loop reduction (+:c)
+//  for (int i = 0; i < 10; i++)
+//    {
+//      c = 1;
+//    }
+
+#pragma acc data if (1) copy (b)
+  {
+    #pragma acc kernels
+    {
+      b = a;
+    }
+  }
+
+#pragma acc enter data copyin (b)
+#pragma acc kernels present (b)
+    {
+      b = a;
+    }
+  return b;
+}
+
+int
+main ()
+{
+  int b = oacc_parallel_copy<int> (5);
+  int c = oacc_kernels_copy<int> (5);
+
+  return b + c;
+}
diff --git a/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90 b/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90
new file mode 100644
index 0000000..6977525
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90
@@ -0,0 +1,155 @@
+! Exercise combined OpenACC directives.
+
+! { dg-do compile }
+! { dg-options "-fopenacc -fdump-tree-gimple" }
+
+! This error is temporary.  Remove when support is added for these clauses
+! in the middle end.
+! { dg-prune-output "sorry, unimplemented" }
+
+! Update the reduction tests.
+
+subroutine test
+  implicit none
+  integer a(100), i, j, z
+
+  ! PARALLEL
+  
+  !$acc parallel loop collapse (2)
+  do i = 1, 100
+     do j = 1, 10
+     end do
+  end do
+  !$acc end parallel loop
+  
+  !$acc parallel loop gang
+  do i = 1, 100
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop worker
+  do i = 1, 100
+     do j = 1, 10
+     end do
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop vector
+  do i = 1, 100
+     do j = 1, 10
+     end do
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop seq
+  do i = 1, 100
+     do j = 1, 10
+     end do
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop auto
+  do i = 1, 100
+     do j = 1, 10
+     end do
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop tile (2, 3)
+  do i = 1, 100
+     do j = 1, 10
+     end do
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop independent
+  do i = 1, 100
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel loop private (z)
+  do i = 1, 100
+     z = 0
+  end do
+  !$acc end parallel loop
+
+!  !$acc parallel loop reduction (+:z) copy (z)
+!  do i = 1, 100
+!  end do
+!  !$acc end parallel loop
+
+  ! KERNELS
+
+  !$acc kernels loop collapse (2)
+  do i = 1, 100
+     do j = 1, 10
+     end do
+  end do
+  !$acc end kernels loop
+  
+  !$acc kernels loop gang
+  do i = 1, 100
+  end do
+  !$acc end kernels loop
+
+  !$acc kernels loop worker
+  do i = 1, 100
+     do j = 1, 10
+     end do
+  end do
+  !$acc end kernels loop
+
+  !$acc kernels loop vector
+  do i = 1, 100
+     do j = 1, 10
+     end do
+  end do
+  !$acc end kernels loop
+
+  !$acc kernels loop seq
+  do i = 1, 100
+     do j = 1, 10
+     end do
+  end do
+  !$acc end kernels loop
+
+  !$acc kernels loop auto
+  do i = 1, 100
+     do j = 1, 10
+     end do
+  end do
+  !$acc end kernels loop
+
+  !$acc kernels loop tile (2, 3)
+  do i = 1, 100
+     do j = 1, 10
+     end do
+  end do
+  !$acc end kernels loop
+
+  !$acc kernels loop independent
+  do i = 1, 100
+  end do
+  !$acc end kernels loop
+
+  !$acc kernels loop private (z)
+  do i = 1, 100
+     z = 0
+  end do
+  !$acc end kernels loop
+
+!  !$acc kernels loop reduction (+:z) copy (z)
+!  do i = 1, 100
+!  end do
+!  !$acc end kernels loop
+end subroutine test
+
+! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. collapse.2." 2 "gimple" } }
+! { dg-final { scan-tree-dump-times "acc loop private.i. gang" 2 "gimple" } }
+! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. worker" 2 "gimple" } }
+! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. vector" 2 "gimple" } }
+! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. seq" 2 "gimple" } }
+! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. auto" 2 "gimple" } }
+! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. tile.2, 3" 2 "gimple" } }
+! { dg-final { scan-tree-dump-times "acc loop private.i. independent" 2 "gimple" } }
+! { dg-final { scan-tree-dump-times "private.z" 2 "gimple" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/default.f95 b/gcc/testsuite/gfortran.dg/goacc/default.f95
new file mode 100644
index 0000000..c1fc52e
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/default.f95
@@ -0,0 +1,17 @@
+! { dg-do compile }
+
+program tile
+  integer i, j, a
+
+  !$acc parallel default (shared) ! { dg-error "Unclassifiable OpenACC directive" }
+  !$acc end parallel ! { dg-error "Unexpected" }
+
+  !$acc parallel default (private) ! { dg-error "Unclassifiable OpenACC directive" }
+  !$acc end parallel ! { dg-error "Unexpected" }
+
+  !$acc parallel default (none)
+  !$acc end parallel
+
+  !$acc parallel default (firstprivate) ! { dg-error "Unclassifiable OpenACC directive" }
+  !$acc end parallel ! { dg-error "Unexpected" }
+end program tile
diff --git a/gcc/testsuite/gfortran.dg/goacc/default_none.f95 b/gcc/testsuite/gfortran.dg/goacc/default_none.f95
new file mode 100644
index 0000000..5ce66ae
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/default_none.f95
@@ -0,0 +1,59 @@
+! Ensure that the internal array variables, offset, lbound, etc., don't
+! trigger errors with default(none).
+
+! { dg-do compile }
+
+program main
+  implicit none
+  integer i
+  integer,parameter :: n = 100
+  integer,allocatable :: a1(:), a2(:,:)
+
+  allocate (a1 (n))
+  allocate (a2 (-n:n,-n:n))
+  a1 (:) = -1
+
+  !$acc parallel loop default(none) copy (a1(1:n))
+  do i = 1,n
+     a1(i) = i
+  end do
+  !$acc end parallel loop
+
+  call foo (a1)
+  call bar (a1, n)
+  call foobar (a2,n)
+
+contains
+
+  subroutine foo (da1)
+    integer :: da1(n)
+
+    !$acc parallel loop default(none) copy (da1(1:n))
+    do i = 1,n
+       da1(i) = i*2
+    end do
+    !$acc end parallel loop
+  end subroutine foo
+end program main
+
+subroutine bar (da2,n)
+  integer :: n, da2(n)
+  integer i
+
+  !$acc parallel loop default(none) copy (da2(1:n)) firstprivate(n)
+  do i = 1,n
+     da2(i) = i*3
+  end do
+  !$acc end parallel loop
+end subroutine bar
+
+subroutine foobar (da3,n)
+  integer :: n, da3(-n:n,-n:n)
+  integer i
+
+  !$acc parallel loop default(none) copy (da3(-n:n,-n:n)) firstprivate(n)
+  do i = 1,n
+     da3(i,0) = i*3
+  end do
+  !$acc end parallel loop
+end subroutine foobar
diff --git a/gcc/testsuite/gfortran.dg/goacc/firstprivate-1.f95 b/gcc/testsuite/gfortran.dg/goacc/firstprivate-1.f95
new file mode 100644
index 0000000..fb92dee
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/firstprivate-1.f95
@@ -0,0 +1,11 @@
+! { dg-do compile }
+
+program test
+  integer a, b(100)
+
+  !$acc parallel firstprivate (a, b)
+  !$acc end parallel
+
+  !$acc parallel firstprivate (b(10:20)) ! { dg-error "Syntax error in OpenMP variable list" }
+  !$acc end parallel ! { dg-error "Unexpected !\\\$ACC END PARALLEL statement" }
+end program test
diff --git a/gcc/testsuite/gfortran.dg/goacc/gang-static.f95 b/gcc/testsuite/gfortran.dg/goacc/gang-static.f95
new file mode 100644
index 0000000..4e46cf3
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/gang-static.f95
@@ -0,0 +1,68 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-omplower" }
+
+program main
+  integer, parameter :: n = 100
+  integer i, a(n), b(n)
+
+  do i = 1, n
+     b(i) = i
+  end do
+
+  !$acc parallel loop gang (static:*) num_gangs (10)
+  do i = 1, n
+     a(i) = b(i) + 0
+  end do
+  !$acc end parallel loop
+
+  call test (a, b, 0, n)
+
+  !$acc parallel loop gang (static:1) num_gangs (10)
+  do i = 1, n
+     a(i) = b(i) + 1
+  end do
+  !$acc end parallel loop
+
+  call test (a, b, 1, n)
+
+  !$acc parallel loop gang (static:2) num_gangs (10)
+  do i = 1, n
+     a(i) = b(i) + 2
+  end do
+  !$acc end parallel loop
+
+  call test (a, b, 2, n)
+
+  !$acc parallel loop gang (static:5) num_gangs (10)
+  do i = 1, n
+     a(i) = b(i) + 5
+  end do
+  !$acc end parallel loop
+
+  call test (a, b, 5, n)
+
+  !$acc parallel loop gang (static:20) num_gangs (10)
+  do i = 1, n
+     a(i) = b(i) + 20
+  end do
+  !$acc end parallel loop
+
+  call test (a, b, 20, n)
+
+end program main
+
+subroutine test (a, b, sarg, n)
+  integer n
+  integer a (n), b(n), sarg
+  integer i
+
+  do i = 1, n
+     if (a(i) .ne. b(i) + sarg) call abort ()
+  end do
+end subroutine test
+
+! { dg-final { scan-tree-dump-times "gang\\(static:\\\*\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "gang\\(static:1\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "gang\\(static:2\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "gang\\(static:5\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "gang\\(static:20\\)" 1 "omplower" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-inner.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-inner.f95
new file mode 100644
index 0000000..4db3a50
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-inner.f95
@@ -0,0 +1,23 @@
+! { dg-additional-options "-O2" }
+! { dg-additional-options "-ftree-parallelize-loops=32" }
+
+program main
+   implicit none
+
+   integer :: a(100,100), b(100,100)
+   integer :: i, j, d
+
+   !$acc kernels
+   do i=1,100
+     do j=1,100
+       a(i,j) = 1
+       b(i,j) = 2
+       a(i,j) = a(i,j) + b(i,j)
+     end do
+   end do
+   !$acc end kernels
+
+   d = sum(a)
+
+   print *,d
+end program main
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loops-adjacent.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loops-adjacent.f95
new file mode 100644
index 0000000..fef3d10
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loops-adjacent.f95
@@ -0,0 +1,19 @@
+! { dg-additional-options "-O2" }
+! { dg-additional-options "-ftree-parallelize-loops=10" }
+
+program main
+   implicit none
+
+   integer :: a(10000), b(10000)
+   integer :: d
+
+   !$acc kernels
+   a = 1
+   b = 2
+   a = a + b
+   !$acc end kernels
+
+   d = sum(a)
+
+   print *,d
+end program main
diff --git a/gcc/testsuite/gfortran.dg/goacc/list.f95 b/gcc/testsuite/gfortran.dg/goacc/list.f95
index 94fdadd..a8006bc 100644
--- a/gcc/testsuite/gfortran.dg/goacc/list.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/list.f95
@@ -5,7 +5,7 @@ program test
   implicit none
 
   integer :: i, j, k, l, a(10)
-  common /b/ j, k
+  common /b/ k
   real, pointer :: p1 => NULL()
   complex :: c, d(10)
 
@@ -108,4 +108,4 @@ program test
   !$acc host_data use_device(p1) ! { dg-error "POINTER" }
   !$acc end host_data
 
-end program test
\ No newline at end of file
+end program test
diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-2.f95 b/gcc/testsuite/gfortran.dg/goacc/loop-2.f95
index f85691e..b5e6368 100644
--- a/gcc/testsuite/gfortran.dg/goacc/loop-2.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/loop-2.f95
@@ -66,7 +66,7 @@ program test
     !$acc loop seq worker ! { dg-error "conflicts with" }
     DO i = 1,10
     ENDDO
-    !$acc loop gang worker ! { dg-error "conflicts with" }
+    !$acc loop gang worker
     DO i = 1,10
     ENDDO
 
@@ -94,10 +94,10 @@ program test
     !$acc loop seq vector ! { dg-error "conflicts with" }
     DO i = 1,10
     ENDDO
-    !$acc loop gang vector ! { dg-error "conflicts with" }
+    !$acc loop gang vector
     DO i = 1,10
     ENDDO
-    !$acc loop worker vector ! { dg-error "conflicts with" }
+    !$acc loop worker vector
     DO i = 1,10
     ENDDO
 
@@ -239,7 +239,7 @@ program test
     !$acc loop seq worker ! { dg-error "conflicts with" }
     DO i = 1,10
     ENDDO
-    !$acc loop gang worker ! { dg-error "conflicts with" }
+    !$acc loop gang worker
     DO i = 1,10
     ENDDO
 
@@ -267,10 +267,10 @@ program test
     !$acc loop seq vector ! { dg-error "conflicts with" }
     DO i = 1,10
     ENDDO
-    !$acc loop gang vector ! { dg-error "conflicts with" }
+    !$acc loop gang vector
     DO i = 1,10
     ENDDO
-    !$acc loop worker vector ! { dg-error "conflicts with" }
+    !$acc loop worker vector
     DO i = 1,10
     ENDDO
 
@@ -392,7 +392,7 @@ program test
   !$acc kernels loop seq worker ! { dg-error "conflicts with" }
   DO i = 1,10
   ENDDO
-  !$acc kernels loop gang worker ! { dg-error "conflicts with" }
+  !$acc kernels loop gang worker
   DO i = 1,10
   ENDDO
 
@@ -420,10 +420,10 @@ program test
   !$acc kernels loop seq vector ! { dg-error "conflicts with" }
   DO i = 1,10
   ENDDO
-  !$acc kernels loop gang vector ! { dg-error "conflicts with" }
+  !$acc kernels loop gang vector
   DO i = 1,10
   ENDDO
-  !$acc kernels loop worker vector ! { dg-error "conflicts with" }
+  !$acc kernels loop worker vector
   DO i = 1,10
   ENDDO
 
@@ -544,7 +544,7 @@ program test
   !$acc parallel loop seq worker ! { dg-error "conflicts with" }
   DO i = 1,10
   ENDDO
-  !$acc parallel loop gang worker ! { dg-error "conflicts with" }
+  !$acc parallel loop gang worker
   DO i = 1,10
   ENDDO
 
@@ -572,10 +572,10 @@ program test
   !$acc parallel loop seq vector ! { dg-error "conflicts with" }
   DO i = 1,10
   ENDDO
-  !$acc parallel loop gang vector ! { dg-error "conflicts with" }
+  !$acc parallel loop gang vector
   DO i = 1,10
   ENDDO
-  !$acc parallel loop worker vector ! { dg-error "conflicts with" }
+  !$acc parallel loop worker vector
   DO i = 1,10
   ENDDO
 
@@ -646,4 +646,4 @@ program test
   !$acc parallel loop gang worker tile(*) 
   DO i = 1,10
   ENDDO
-end
\ No newline at end of file
+end
diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-4.f95 b/gcc/testsuite/gfortran.dg/goacc/loop-4.f95
new file mode 100644
index 0000000..7c53c02
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/loop-4.f95
@@ -0,0 +1,7 @@
+! Ensure that loops not affiliated with acc compute regions cause an error.
+
+subroutine test1
+    !$acc loop gang ! { dg-error "loop directive must be associated with an OpenACC compute region" }
+  DO i = 1,10
+  ENDDO
+end subroutine test1
diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-5.f95 b/gcc/testsuite/gfortran.dg/goacc/loop-5.f95
new file mode 100644
index 0000000..5cbd975
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/loop-5.f95
@@ -0,0 +1,363 @@
+! { dg-do compile }
+! { dg-additional-options "-fmax-errors=100" }
+
+! { dg-prune-output "sorry, unimplemented" }
+! { dg-prune-output "Error: work-sharing region" }
+
+program test
+  implicit none
+  integer :: i, j
+
+  !$acc kernels
+    !$acc loop auto
+    DO i = 1,10
+    ENDDO
+    !$acc loop gang
+    DO i = 1,10
+    ENDDO
+    !$acc loop gang(5)
+    DO i = 1,10
+    ENDDO
+    !$acc loop gang(num:5)
+    DO i = 1,10
+    ENDDO
+    !$acc loop gang(static:5)
+    DO i = 1,10
+    ENDDO
+    !$acc loop gang(static:*)
+    DO i = 1,10
+    ENDDO
+    !$acc loop gang
+    DO i = 1,10
+      !$acc loop vector
+      DO j = 1,10
+      ENDDO
+      !$acc loop worker
+      DO j = 1,10
+      ENDDO
+    ENDDO
+
+    !$acc loop worker
+    DO i = 1,10
+    ENDDO
+    !$acc loop worker(5)
+    DO i = 1,10
+    ENDDO
+    !$acc loop worker(num:5)
+    DO i = 1,10
+    ENDDO
+    !$acc loop worker
+    DO i = 1,10
+      !$acc loop vector
+      DO j = 1,10
+      ENDDO
+    ENDDO
+    !$acc loop gang worker
+    DO i = 1,10
+    ENDDO
+
+    !$acc loop vector
+    DO i = 1,10
+    ENDDO
+    !$acc loop vector(5)
+    DO i = 1,10
+    ENDDO
+    !$acc loop vector(length:5)
+    DO i = 1,10
+    ENDDO
+    !$acc loop vector
+    DO i = 1,10
+    ENDDO
+    !$acc loop gang vector
+    DO i = 1,10
+    ENDDO
+    !$acc loop worker vector
+    DO i = 1,10
+    ENDDO
+
+    !$acc loop auto
+    DO i = 1,10
+    ENDDO
+
+    !$acc loop tile(1)
+    DO i = 1,10
+    ENDDO
+    !$acc loop tile(2)
+    DO i = 1,10
+    ENDDO
+    !$acc loop tile(6-2)
+    DO i = 1,10
+    ENDDO
+    !$acc loop tile(6+2)
+    DO i = 1,10
+    ENDDO
+    !$acc loop tile(*)
+    DO i = 1,10
+    ENDDO
+    !$acc loop tile(*, 1)
+    DO i = 1,10
+      DO j = 1,10
+      ENDDO
+    ENDDO
+    !$acc loop tile(-1) ! { dg-warning "must be positive" }
+    do i = 1,10
+    enddo
+    !$acc loop vector tile(*)
+    DO i = 1,10
+    ENDDO
+    !$acc loop worker tile(*)
+    DO i = 1,10
+    ENDDO
+    !$acc loop gang tile(*)
+    DO i = 1,10
+    ENDDO
+    !$acc loop vector gang tile(*)
+    DO i = 1,10
+    ENDDO
+    !$acc loop vector worker tile(*)
+    DO i = 1,10
+    ENDDO
+    !$acc loop gang worker tile(*)
+    DO i = 1,10
+    ENDDO
+  !$acc end kernels
+
+
+  !$acc parallel
+    !$acc loop tile(1)
+    DO i = 1,10
+    ENDDO
+    !$acc loop tile(*)
+    DO i = 1,10
+    ENDDO
+    !$acc loop tile(2)
+    DO i = 1,10
+      DO j = 1,10
+      ENDDO
+    ENDDO
+    !$acc loop tile(-1) ! { dg-warning "must be positive" }
+    do i = 1,10
+    enddo
+    !$acc loop vector tile(*)
+    DO i = 1,10
+    ENDDO
+    !$acc loop worker tile(*)
+    DO i = 1,10
+    ENDDO
+    !$acc loop gang tile(*)
+    DO i = 1,10
+    ENDDO
+    !$acc loop vector gang tile(*)
+    DO i = 1,10
+    ENDDO
+    !$acc loop vector worker tile(*)
+    DO i = 1,10
+    ENDDO
+    !$acc loop gang worker tile(*)
+    DO i = 1,10
+    ENDDO
+  !$acc end parallel
+
+  !$acc kernels loop auto
+  DO i = 1,10
+  ENDDO
+  !$acc kernels loop gang
+  DO i = 1,10
+  ENDDO
+  !$acc kernels loop gang(5)
+  DO i = 1,10
+  ENDDO
+  !$acc kernels loop gang(num:5)
+  DO i = 1,10
+  ENDDO
+  !$acc kernels loop gang(static:5)
+  DO i = 1,10
+  ENDDO
+  !$acc kernels loop gang(static:*)
+  DO i = 1,10
+  ENDDO
+  !$acc kernels loop gang
+  DO i = 1,10
+    !$acc kernels loop gang ! { dg-error "OpenACC construct inside of non-OpenACC region" }
+    DO j = 1,10
+    ENDDO
+  ENDDO
+
+  !$acc kernels loop worker
+  DO i = 1,10
+  ENDDO
+  !$acc kernels loop worker(5)
+  DO i = 1,10
+  ENDDO
+  !$acc kernels loop worker(num:5)
+  DO i = 1,10
+  ENDDO
+  !$acc kernels loop worker
+  DO i = 1,10
+    !$acc kernels loop worker ! { dg-error "OpenACC construct inside of non-OpenACC region" }
+    DO j = 1,10
+    ENDDO
+    !$acc kernels loop gang ! { dg-error "OpenACC construct inside of non-OpenACC region" }
+    DO j = 1,10
+    ENDDO
+  ENDDO
+  !$acc kernels loop gang worker
+  DO i = 1,10
+  ENDDO
+
+  !$acc kernels loop vector
+  DO i = 1,10
+  ENDDO
+  !$acc kernels loop vector(5)
+  DO i = 1,10
+  ENDDO
+  !$acc kernels loop vector(length:5)
+  DO i = 1,10
+  ENDDO
+  !$acc kernels loop vector
+  DO i = 1,10
+    !$acc kernels loop vector ! { dg-error "OpenACC construct inside of non-OpenACC region" }
+    DO j = 1,10
+    ENDDO
+    !$acc kernels loop worker ! { dg-error "OpenACC construct inside of non-OpenACC region" }
+    DO j = 1,10
+    ENDDO
+    !$acc kernels loop gang ! { dg-error "OpenACC construct inside of non-OpenACC region" }
+    DO j = 1,10
+    ENDDO
+  ENDDO
+  !$acc kernels loop gang vector
+  DO i = 1,10
+  ENDDO
+  !$acc kernels loop worker vector
+  DO i = 1,10
+  ENDDO
+
+  !$acc kernels loop auto
+  DO i = 1,10
+  ENDDO
+
+  !$acc kernels loop tile(1)
+  DO i = 1,10
+  ENDDO
+  !$acc kernels loop tile(*)
+  DO i = 1,10
+  ENDDO
+  !$acc kernels loop tile(*, 1)
+  DO i = 1,10
+    DO j = 1,10
+    ENDDO
+  ENDDO
+  !$acc kernels loop tile(-1) ! { dg-warning "must be positive" }
+  do i = 1,10
+  enddo
+  !$acc kernels loop vector tile(*)
+  DO i = 1,10
+  ENDDO
+  !$acc kernels loop worker tile(*)
+  DO i = 1,10
+  ENDDO
+  !$acc kernels loop gang tile(*)
+  DO i = 1,10
+  ENDDO
+  !$acc kernels loop vector gang tile(*)
+  DO i = 1,10
+  ENDDO
+  !$acc kernels loop vector worker tile(*)
+  DO i = 1,10
+  ENDDO
+  !$acc kernels loop gang worker tile(*)
+  DO i = 1,10
+  ENDDO
+
+  !$acc parallel loop auto
+  DO i = 1,10
+  ENDDO
+  !$acc parallel loop gang
+  DO i = 1,10
+  ENDDO
+  !$acc parallel loop gang(static:5)
+  DO i = 1,10
+  ENDDO
+  !$acc parallel loop gang(static:*)
+  DO i = 1,10
+  ENDDO
+  !$acc parallel loop gang
+  DO i = 1,10
+    !$acc parallel loop gang ! { dg-error "OpenACC construct inside of non-OpenACC region" }
+    DO j = 1,10
+    ENDDO
+  ENDDO
+
+  !$acc parallel loop worker
+  DO i = 1,10
+  ENDDO
+  !$acc parallel loop worker
+  DO i = 1,10
+    !$acc parallel loop worker ! { dg-error "OpenACC construct inside of non-OpenACC region" }
+    DO j = 1,10
+    ENDDO
+    !$acc parallel loop gang ! { dg-error "OpenACC construct inside of non-OpenACC region" }
+    DO j = 1,10
+    ENDDO
+  ENDDO
+  !$acc parallel loop gang worker
+  DO i = 1,10
+  ENDDO
+
+  !$acc parallel loop vector
+  DO i = 1,10
+    !$acc parallel loop vector ! { dg-error "OpenACC construct inside of non-OpenACC region" }
+    DO j = 1,10
+    ENDDO
+    !$acc parallel loop worker ! { dg-error "OpenACC construct inside of non-OpenACC region" }
+    DO j = 1,10
+    ENDDO
+    !$acc parallel loop gang ! { dg-error "OpenACC construct inside of non-OpenACC region" }
+    DO j = 1,10
+    ENDDO
+  ENDDO
+  !$acc parallel loop gang vector
+  DO i = 1,10
+  ENDDO
+  !$acc parallel loop worker vector
+  DO i = 1,10
+  ENDDO
+
+  !$acc parallel loop auto
+  DO i = 1,10
+  ENDDO
+
+  !$acc parallel loop tile(1)
+  DO i = 1,10
+  ENDDO
+  !$acc parallel loop tile(*)
+  DO i = 1,10
+  ENDDO
+  !$acc parallel loop tile(*, 1)
+  DO i = 1,10
+    DO j = 1,10
+    ENDDO
+  ENDDO
+  !$acc parallel loop tile(-1) ! { dg-warning "must be positive" }
+  do i = 1,10
+  enddo
+  !$acc parallel loop vector tile(*)
+  DO i = 1,10
+  ENDDO
+  !$acc parallel loop worker tile(*)
+  DO i = 1,10
+  ENDDO
+  !$acc parallel loop gang tile(*)
+  DO i = 1,10
+  ENDDO
+  !$acc parallel loop vector gang tile(*)
+  DO i = 1,10
+  ENDDO
+  !$acc parallel loop vector worker tile(*)
+  DO i = 1,10
+  ENDDO
+  !$acc parallel loop gang worker tile(*)
+  DO i = 1,10
+  ENDDO
+end
diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-6.f95 b/gcc/testsuite/gfortran.dg/goacc/loop-6.f95
new file mode 100644
index 0000000..e13abc7
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/loop-6.f95
@@ -0,0 +1,82 @@
+! { dg-do compile }
+! { dg-additional-options "-fmax-errors=100" }
+
+! This error is temporary.  Remove when support is added for these clauses
+! in the middle end.
+! { dg-prune-output "sorry, unimplemented" }
+! { dg-prune-output "Error: work-sharing region" }
+
+program test
+  implicit none
+  integer :: i, j
+
+  !$acc parallel
+    !$acc loop auto
+    DO i = 1,10
+    ENDDO
+    !$acc loop gang
+    DO i = 1,10
+    ENDDO
+    !$acc loop gang(static:5)
+    DO i = 1,10
+    ENDDO
+    !$acc loop gang(static:*)
+    DO i = 1,10
+    ENDDO
+    !$acc loop gang
+    DO i = 1,10
+      !$acc loop vector
+      DO j = 1,10
+      ENDDO
+      !$acc loop worker
+      DO j = 1,10
+      ENDDO
+    ENDDO
+
+    !$acc loop worker
+    DO i = 1,10
+    ENDDO
+    !$acc loop worker
+    DO i = 1,10
+      !$acc loop vector
+      DO j = 1,10
+      ENDDO
+    ENDDO
+    !$acc loop gang worker
+    DO i = 1,10
+    ENDDO
+
+    !$acc loop vector
+    DO i = 1,10
+    ENDDO
+    !$acc loop vector(5) ! { dg-error "argument not permitted" }
+    DO i = 1,10
+    ENDDO
+    !$acc loop vector(length:5) ! { dg-error "argument not permitted" }
+    DO i = 1,10
+    ENDDO
+    !$acc loop vector
+    DO i = 1,10
+    ENDDO
+    !$acc loop gang vector
+    DO i = 1,10
+    ENDDO
+    !$acc loop worker vector
+    DO i = 1,10
+    ENDDO
+
+    !$acc loop auto
+    DO i = 1,10
+    ENDDO
+  !$acc end parallel
+
+  !$acc parallel loop vector
+  DO i = 1,10
+  ENDDO
+  !$acc parallel loop vector(5) ! { dg-error "argument not permitted" }
+  DO i = 1,10
+  ENDDO
+  !$acc parallel loop vector(length:5) ! { dg-error "argument not permitted" }
+  DO i = 1,10
+  ENDDO
+end
diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90 b/gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90
index d72bae4..6cfd715 100644
--- a/gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90
@@ -3,6 +3,11 @@
 
 ! test for tree-dump-original and spaces-commas
 
+! This error is temporary.  Remove when support is added for these clauses
+! in the middle end.
+! { dg-prune-output "sorry, unimplemented" }
+! { dg-prune-output "Error: work-sharing region" }
+
 program test
   implicit none
   integer :: i, j, k, m, sum
@@ -17,7 +22,7 @@ program test
 
   !$acc loop independent gang (3)
   DO i = 1,10
-    !$acc loop worker(3) ! { dg-error "work-sharing region may not be closely nested inside of work-sharing, critical, ordered, master or explicit task region" }
+    !$acc loop worker(3)
     DO j = 1,10
       !$acc loop vector(5)
       DO k = 1,10
diff --git a/gcc/testsuite/gfortran.dg/goacc/multi-clause.f90 b/gcc/testsuite/gfortran.dg/goacc/multi-clause.f90
new file mode 100644
index 0000000..2870076
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/multi-clause.f90
@@ -0,0 +1,13 @@
+! Test if variable appearing in multiple clauses are errors.
+
+! { dg-compile }
+
+program combined
+  implicit none
+  integer a(100), i, j
+
+  !$acc parallel loop reduction (+:j) copy (j) copyout(j) ! { dg-error "Symbol 'j' present on multiple clauses" }
+  do i = 1, 100
+  end do
+  !$acc end parallel loop
+end program combined
diff --git a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95
index 4915744..9037f6c 100644
--- a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95
@@ -37,4 +37,3 @@ end program test
 
 ! { dg-final { scan-tree-dump-times "map\\(force_deviceptr:u\\)" 1 "original" } } 
 ! { dg-final { scan-tree-dump-times "private\\(v\\)" 1 "original" } } 
-! { dg-final { scan-tree-dump-times "firstprivate\\(w\\)" 1 "original" } } 
diff --git a/gcc/testsuite/gfortran.dg/goacc/update.f95 b/gcc/testsuite/gfortran.dg/goacc/update.f95
new file mode 100644
index 0000000..d88d20e
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/update.f95
@@ -0,0 +1,5 @@
+! { dg-do compile } 
+
+program foo
+  !$acc update ! { dg-error "must contain at least one 'device' or 'host' or 'self' clause" }
+end program foo
diff --git a/gcc/tree-core.h b/gcc/tree-core.h
index d913e51..f2299f2 100644
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -432,7 +432,10 @@ enum omp_clause_code {
   OMP_CLAUSE_NUM_WORKERS,
 
   /* OpenACC clause: vector_length (integer-expression).  */
-  OMP_CLAUSE_VECTOR_LENGTH
+  OMP_CLAUSE_VECTOR_LENGTH,
+
+  /* OpenACC clause: tile ( size-expr-list ).  */
+  OMP_CLAUSE_TILE
 };
 
 #undef DEFTREESTRUCT
diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index efae4c0..b1685fd 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -932,6 +932,12 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, int flags)
     case OMP_CLAUSE_INDEPENDENT:
       pp_string (pp, "independent");
       break;
+    case OMP_CLAUSE_TILE:
+      pp_string (pp, "tile(");
+      dump_generic_node (pp, OMP_CLAUSE_TILE_LIST (clause),
+			 spc, flags, false);
+      pp_right_paren (pp);
+      break;
 
     default:
       /* Should never happen.  */
diff --git a/gcc/tree.c b/gcc/tree.c
index 18d6544..5b9a7bd 100644
--- a/gcc/tree.c
+++ b/gcc/tree.c
@@ -328,6 +328,7 @@ unsigned const char omp_clause_num_ops[] =
   1, /* OMP_CLAUSE_NUM_GANGS  */
   1, /* OMP_CLAUSE_NUM_WORKERS  */
   1, /* OMP_CLAUSE_VECTOR_LENGTH  */
+  1, /* OMP_CLAUSE_TILE  */
 };
 
 const char * const omp_clause_code_name[] =
@@ -398,7 +399,8 @@ const char * const omp_clause_code_name[] =
   "vector",
   "num_gangs",
   "num_workers",
-  "vector_length"
+  "vector_length",
+  "tile"
 };
 
 
@@ -11595,6 +11597,7 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data,
 	case OMP_CLAUSE_DEFAULTMAP:
 	case OMP_CLAUSE_AUTO:
 	case OMP_CLAUSE_SEQ:
+	case OMP_CLAUSE_TILE:
 	  WALK_SUBTREE_TAIL (OMP_CLAUSE_CHAIN (*tp));
 
 	case OMP_CLAUSE_LASTPRIVATE:
diff --git a/gcc/tree.h b/gcc/tree.h
index dc592d4..6768b3b 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1556,6 +1556,9 @@ extern void protected_set_expr_location (tree, location_t);
 #define OMP_CLAUSE_DEFAULT_KIND(NODE) \
   (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_DEFAULT)->omp_clause.subcode.default_kind)
 
+#define OMP_CLAUSE_TILE_LIST(NODE) \
+  OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_TILE), 0)
+
 /* SSA_NAME accessors.  */
 
 /* Returns the IDENTIFIER_NODE giving the SSA name a name or NULL_TREE