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

3 files changed, 139 insertions, 9 deletions

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;