public inbox for gcc-patches@gcc.gnu.org
 help / color / mirror / Atom feed
* [openacc] tile, independent, default, private and firstprivate support in c/++
@ 2015-11-03 22:17 Cesar Philippidis
  2015-11-04 10:24 ` Jakub Jelinek
  2015-11-05 12:15 ` Thomas Schwinge
  0 siblings, 2 replies; 26+ messages in thread
From: Cesar Philippidis @ 2015-11-03 22:17 UTC (permalink / raw)
  To: gcc-patches, Jakub Jelinek

[-- Attachment #1: Type: text/plain, Size: 835 bytes --]

This patch does the following to the c and c++ front ends:

 * parsing support for the tile, independent, default (none),
   private and firstprivate clauses in c and c++

 * updates c_oacc_split_loop_clauses to filter out the loop clauses
   from combined parallel/kernels loops

The c front end already had some support for private and firstprivate in
openacc. However, the c++ front end wasn't associating any of those
clauses with parallel, kernels or acc loops.

For reference, here's the grammar for the tile clause from section 2.7
in version 2.0a of the openacc spec:

    tile( size-expr-list )

  where size-expr is one of:

    *
    int-expr

That '*' symbol complicated the parsing a little, since it's no longer a
primary expression.

I've bootstrapped and regression tested this on x86_64. Is this ok for
trunk?

Cesar


[-- Attachment #2: misc-c-cxx-clauses.diff --]
[-- Type: text/x-patch, Size: 33881 bytes --]

2015-11-03  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_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_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.

diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c
index 133d079..a04caf3 100644
--- a/gcc/c-family/c-omp.c
+++ b/gcc/c-family/c-omp.c
@@ -709,12 +709,21 @@ 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..ab2bd15 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 %<)%>");
 
@@ -11376,6 +11386,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;
+  vec<tree, va_gc> *tvec = make_tree_vector ();
+
+  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 %<(%>"))
+    {
+      release_tree_vector (tvec);
+      return list;
+    }
+
+  do
+    {
+      if (c_parser_next_token_is (parser, CPP_MULT))
+	{
+	  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 +12663,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 +12692,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 +12742,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 +12827,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 +13240,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 +13293,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 +13309,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 a90bf3b..200bf0c 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -29120,6 +29120,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;
@@ -29214,6 +29216,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;
@@ -29714,6 +29718,58 @@ 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, tree list, location_t here)
+{
+  tree c, expr = error_mark_node;
+  location_t loc;
+  tree tile = NULL_TREE;
+  vec<tree, va_gc> *tvec = make_tree_vector ();
+
+  check_no_duplicate_clause (list, OMP_CLAUSE_TILE, "tile", here);
+
+  loc = cp_lexer_peek_token (parser->lexer)->location;
+  if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
+    goto cleanup_error;
+
+  do
+    {
+      if (cp_lexer_next_token_is (parser->lexer, CPP_MULT))
+	{
+	  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 (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 ) */
 
@@ -29860,10 +29916,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;
@@ -29884,7 +29944,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;
@@ -29898,7 +29958,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))
@@ -31467,6 +31530,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";
@@ -31475,6 +31542,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,
@@ -31488,6 +31560,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";
@@ -31516,6 +31594,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";
@@ -31529,6 +31612,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, clauses, here);
+	  c_name = "tile";
+	  break;
 	case PRAGMA_OACC_CLAUSE_VECTOR:
 	  c_name = "vector";
 	  clauses = cp_parser_oacc_shape_clause (parser, OMP_CLAUSE_VECTOR,
@@ -31616,7 +31703,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:
@@ -34460,12 +34547,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,
@@ -34510,6 +34600,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)		\
@@ -34525,7 +34616,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)		\
@@ -34534,6 +34627,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..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;
 
 	case OMP_CLAUSE_DEFAULTMAP:
@@ -7482,6 +7483,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/tile.c b/gcc/testsuite/c-c++-common/goacc/tile.c
new file mode 100644
index 0000000..ce1391a
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/tile.c
@@ -0,0 +1,61 @@
+/* { dg-do compile } */
+
+int
+main ()
+{
+  int i;
+
+#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++)
+    ;
+
+  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

^ permalink raw reply	[flat|nested] 26+ messages in thread

end of thread, other threads:[~2021-07-21 22:06 UTC | newest]

Thread overview: 26+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2015-11-03 22:17 [openacc] tile, independent, default, private and firstprivate support in c/++ Cesar Philippidis
2015-11-04 10:24 ` Jakub Jelinek
2015-11-04 17:55   ` Cesar Philippidis
2015-11-05 16:31     ` Jakub Jelinek
2015-11-05  4:58   ` Cesar Philippidis
2015-11-05  7:29     ` Jakub Jelinek
2015-11-05 14:58       ` Cesar Philippidis
2015-11-05 12:15 ` Thomas Schwinge
2015-11-05 14:48   ` Cesar Philippidis
2015-11-05 17:02     ` Thomas Schwinge
2015-11-05 17:13       ` Nathan Sidwell
2015-11-05 17:27         ` Cesar Philippidis
2015-11-06  2:11         ` Cesar Philippidis
2015-11-06  6:50           ` Jakub Jelinek
2015-11-06 13:44             ` Nathan Sidwell
2015-11-06 13:42           ` Nathan Sidwell
2015-11-09 11:31           ` [gomp4] " Thomas Schwinge
2021-07-21 22:06           ` Thomas Schwinge
     [not found]         ` <563D0735.7070601@mentor.com>
     [not found]           ` <87h9lqaofq.fsf@schwinge.name>
2015-11-06 23:31             ` [gomp4] backport trunk FE changes Cesar Philippidis
2015-11-07 11:45               ` Combined constructs' clause splitting (was: [gomp4] backport trunk FE changes) Thomas Schwinge
2015-11-07 16:14                 ` Combined constructs' clause splitting Cesar Philippidis
2015-11-08 15:46                 ` Tom de Vries
2015-11-19  0:03                   ` Cesar Philippidis
2015-11-19  0:09                     ` Tom de Vries
2015-11-07 12:30               ` [gomp4] backport trunk FE changes Thomas Schwinge
2015-11-07 16:05                 ` Cesar Philippidis

This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).