diff mbox

[openacc] tile, independent, default, private and firstprivate support in c/++

Message ID 563A46A5.6000702@codesourcery.com
State New
Headers show

Commit Message

Cesar Philippidis Nov. 4, 2015, 5:55 p.m. UTC
On 11/04/2015 02:24 AM, Jakub Jelinek wrote:
> On Tue, Nov 03, 2015 at 02:16:59PM -0800, Cesar Philippidis wrote:

>> +
>> +  do
>> +    {
>> +      if (c_parser_next_token_is (parser, CPP_MULT))
>> +	{
>> +	  c_parser_consume_token (parser);
>> +	  expr = integer_minus_one_node;
>> +	}
>> +      else
> 
> Is this right?  If it is either * or (assignment) expression, then
> I'd expect to parse only CPP_MULT followed by CPP_CLOSE_PAREN
> or CPP_COMMA that way (C parser has 2 tokens look-ahead, so it should be
> fine), so that
> tile (a + b + c, *)
> is parsed as
> (a + b + c); -1
> and so is
> tile (*, a + b)
> as
> -1; (a + b)
> while
> tile (*a, *b)
> is
> *a; *b.
> 
> Guess the gang clause parsing that went into the trunk already has the
> same bug,
> gang (static: *)
> or
> gang (static: *, num: 5)
> should be special, but
> gang (static: *ptr)
> should be
> gang (static: (*ptr))

Thanks for catching that. I'll fix the tile and shape scanners in both
front ends.

>> diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
>> index c73dcd0..14d006b 100644
>> --- a/gcc/cp/semantics.c
>> +++ b/gcc/cp/semantics.c
>> @@ -6704,9 +6704,51 @@ 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:
>> +	  {
>> +	    tree list = OMP_CLAUSE_TILE_LIST (c);
>> +
>> +	    while (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;
>> +		list = TREE_CHAIN (list);
>> +	      }
>> +	  }
>> +	  break;
> 
> Have you verified pt.c does the right thing when instantiating the
> OMP_CLAUSE_TILE clause (I mean primarily the TREE_VEC in there)?
> There really should be testcases for that.

I don't think we have any support for templates in trunk yet. Should I
add it to this patch, or should I address that in a follow up patch?

By the way, template support for num_gangs, num_worker and vector_length
depend on this patch
<https://gcc.gnu.org/ml/gcc-patches/2015-10/msg03455.html> because the
c++ front end is currently incorrectly trying to do type checking as it
scans those clauses.

>> diff --git a/gcc/gimplify.c b/gcc/gimplify.c
>> index 03203c0..08b192d 100644
>> --- a/gcc/gimplify.c
>> +++ b/gcc/gimplify.c
>> @@ -6997,7 +6997,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
>>  
>>  	case OMP_CLAUSE_DEVICE_RESIDENT:
>>  	case OMP_CLAUSE_USE_DEVICE:
>> -	case OMP_CLAUSE_INDEPENDENT:
>>  	  remove = true;
>>  	  break;
>>  
>> @@ -7007,6 +7006,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:
>> @@ -7014,6 +7014,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
>>  	case OMP_CLAUSE_NOGROUP:
>>  	case OMP_CLAUSE_THREADS:
>>  	case OMP_CLAUSE_SIMD:
>> +	case OMP_CLAUSE_TILE:
>>  	  break;
> 
> No gimplification of the expressions in the tile clause?

Ultimately we're just ignoring it, but for completeness' sake I fixed this.

I'm still testing this patch, but something like this OK for trunk?

cesar

Comments

Jakub Jelinek Nov. 5, 2015, 4:30 p.m. UTC | #1
Hi!

On Wed, Nov 04, 2015 at 09:55:49AM -0800, Cesar Philippidis wrote:

So, you are going to deal with gang parsing incrementally?
Fine with me.

> +/* 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;
> +
> +  vec<tree, va_gc> *tvec = make_tree_vector ();

Seems I've misread your patch, thought you are using TREE_VEC, while
you are actually using TREE_LIST, but populating it in a weird way.
I think more efficient would be just to
  tree tile = NULL_TREE;
here, then:

> +
> +      vec_safe_push (tvec, expr);

  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 = build_tree_list_vec (tvec);

  tile = nreverse (tile);
  
> +  OMP_CLAUSE_TILE_LIST (c) = tile;
> +  OMP_CLAUSE_CHAIN (c) = list;
> +  release_tree_vector (tvec);

and remove the release_tree_vector calls.

> +static tree
> +cp_parser_oacc_clause_tile (cp_parser *parser, location_t clause_loc, tree list)

This is already too long line.

> +	case OMP_CLAUSE_TILE:
> +	  {
> +	    tree list = OMP_CLAUSE_TILE_LIST (c);
> +
> +	    while (list)

I'd say
	    for (tree list = OMP_CLAUSE_TILE_LIST (c);
		 list; list = TREE_CHAIN (list))
would be more readable.

> --- a/gcc/gimplify.c
> +++ b/gcc/gimplify.c
> @@ -6995,9 +6995,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;
> +	    }

After all, you are already using for here ;)

Otherwise LGTM, but please clear with Thomas, I think he had some issues
with the patch.

	Jakub
diff mbox

Patch

2015-11-04  Cesar Philippidis  <cesar@codesourcery.com>
	    Thomas Schwinge  <thomas@codesourcery.com>
	    James Norris  <jnorris@codesourcery.com>

	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 and independent as loop clauses.  Associate REDUCTION
	clauses with parallel and kernels.
	* c-pragma.h (enum pragma_omp_clause): Add entries for
	PRAGMA_OACC_CLAUSE_{INDEPENDENT,TILE,DEFAULT}.

	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-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}.
	* semantics.c (finish_omp_clauses): Add support for
	OMP_CLAUSE_INDEPENDENT and OMP_CLAUSE_TILE.

	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.

diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c
index ca64eda..1306e09 100644
--- a/gcc/c-family/c-omp.c
+++ b/gcc/c-family/c-omp.c
@@ -709,12 +709,20 @@  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:
 	  OMP_CLAUSE_CHAIN (clauses) = loop_clauses;
 	  loop_clauses = clauses;
 	  break;
 
+	  /* 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/c-parser.c b/gcc/c/c-parser.c
index ec88c65..8360c35 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -10006,6 +10006,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;
@@ -10102,6 +10104,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;
@@ -10539,10 +10543,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;
@@ -10563,7 +10570,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;
@@ -10577,7 +10584,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 %<)%>");
 
@@ -11234,7 +11244,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;
@@ -11376,6 +11389,83 @@  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;
+
+  vec<tree, va_gc> *tvec = make_tree_vector ();
+
+  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)
+	    goto cleanup_error;
+
+	  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;
+	    }
+	}
+
+      vec_safe_push (tvec, expr);
+      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 = build_tree_list_vec (tvec);
+  OMP_CLAUSE_TILE_LIST (c) = tile;
+  OMP_CLAUSE_CHAIN (c) = list;
+  release_tree_vector (tvec);
+  return c;
+
+ cleanup_error:
+  c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, "expected %<)%>");
+  release_tree_vector (tvec);
+  return list;
+}
+
 /* OpenACC:
    wait ( int-expr-list ) */
 
@@ -12576,6 +12666,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";
@@ -12601,6 +12695,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";
@@ -12646,6 +12745,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,
@@ -12727,7 +12830,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:
@@ -13140,13 +13243,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)
@@ -13191,6 +13296,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)		\
@@ -13206,8 +13312,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)		\
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index 2363b9b..24cedf3 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -12946,10 +12946,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/parser.c b/gcc/cp/parser.c
index 12452e6..8f02f4c 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -29119,6 +29119,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;
@@ -29213,6 +29215,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;
@@ -29645,7 +29649,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;
@@ -29713,6 +29720,59 @@  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;
+
+  vec<tree, va_gc> *tvec = make_tree_vector ();
+  
+  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)
+	goto cleanup_error;
+
+      vec_safe_push (tvec, expr);
+
+      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 = build_tree_list_vec (tvec);
+  OMP_CLAUSE_TILE_LIST (c) = tile;
+  OMP_CLAUSE_CHAIN (c) = list;
+  release_tree_vector (tvec);
+  return c;
+
+ cleanup_error:
+  release_tree_vector (tvec);
+  return list;
+}
+
 /* OpenACC:
    vector_length ( expression ) */
 
@@ -29859,10 +29919,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 +29947,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 +29961,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))
@@ -31466,6 +31533,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";
@@ -31474,6 +31545,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,
@@ -31487,6 +31563,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:
 	  clauses = cp_parser_omp_clause_num_gangs (parser, clauses);
 	  c_name = "num_gangs";
@@ -31515,6 +31597,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";
@@ -31528,6 +31615,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,
@@ -31615,7 +31706,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:
@@ -34459,12 +34550,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,
@@ -34509,6 +34603,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)		\
@@ -34524,7 +34619,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)		\
@@ -34533,6 +34630,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) )
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index c73dcd0..14d006b 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -6704,9 +6704,51 @@  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:
+	  {
+	    tree list = OMP_CLAUSE_TILE_LIST (c);
+
+	    while (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;
+		list = TREE_CHAIN (list);
+	      }
+	  }
+	  break;
+
 	case OMP_CLAUSE_INBRANCH:
 	case OMP_CLAUSE_NOTINBRANCH:
 	  if (branch_seen)
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 03203c0..91fde9b 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -6995,9 +6995,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;
 
@@ -7007,6 +7016,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:
@@ -7482,6 +7492,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 ccf0b63..9c87ac9 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -2230,6 +2230,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:
@@ -2408,6 +2409,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/c-c++-common/goacc/combined-directives.c b/gcc/testsuite/c-c++-common/goacc/combined-directives.c
new file mode 100644
index 0000000..cefe712
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/combined-directives.c
@@ -0,0 +1,117 @@ 
+// { dg-do compile }
+// { dg-options "-fopenacc -fdump-tree-gimple" }
+
+// { 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" } }
+// { dg-final { scan-tree-dump-times "reduction..:z. map.force_tofrom:z .len: 4.." 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/tree-core.h b/gcc/tree-core.h
index 6b17da7..562f838 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 8a2398d..6eeb12f 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -905,6 +905,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 65c3117..be18b64a 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1561,6 +1561,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