2015-11-06  Cesar Philippidis  <cesar@codesourcery.com>

	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.

	gcc/c/
	* c-parser.c (c_parser_omp_clause_default): Replace only_none with
	is_oacc argument.
	(c_parser_oacc_shape_clause): Allow pointers arguments to gang static.
	(c_parser_oacc_clause_tile): Backport cleanups from trunnk.
	(c_parser_oacc_all_clauses): Likewise, update call to
	c_parser_omp_clause_default.
	(c_parser_omp_all_clauses): Update call to c_parser_omp_clause_default.

	gcc/cp/
	* parser.c (cp_parser_oacc_shape_clause): Allow pointers arguments to
	gang static.
	(cp_parser_oacc_clause_tile): Backport cleanups from trunnk.
	(cp_parser_omp_clause_default): Replace is_omp argument with is_oacc.
	(cp_parser_oacc_all_clauses): Likewise, update call to
	c_parser_omp_clause_{default,tile}.
	(cp_parser_omp_all_clauses): Update call to
	c_parser_omp_clause_default.
	(OACC_PARALLEL_CLAUSE_MASK): Remove PRAGMA_OACC_CLAUSE_GANG.
	* pt.c (tsubst_omp_clauses):
	* semantics.c (finish_omp_clauses):

	gcc/testsuite/
	* c-c++-common/goacc/combined-directives.c: New test.
	* c-c++-common/goacc/loop-clauses.c: New test.
	* c-c++-common/goacc/loop-shape.c: More test cases.
	* c-c++-common/goacc/loop-tile-k1.c: Update error messages.
	* c-c++-common/goacc/loop-tile-p1.c: Likewise.
	* 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 67d9da0..8411814 100644
--- a/gcc/c-family/c-omp.c
+++ b/gcc/c-family/c-omp.c
@@ -694,13 +694,12 @@ c_finish_omp_for (location_t locus, enum tree_code code, tree declv,
 /* This function splits clauses for OpenACC combined loop
    constructs.  OpenACC combined loop constructs are:
    #pragma acc kernels loop
-   #pragma acc parallel loop
-*/
+   #pragma acc parallel loop  */
 
 tree
 c_oacc_split_loop_clauses (tree clauses, tree *not_loop_clauses)
 {
-  tree next, loop_clauses;
+  tree next, loop_clauses, t;
 
   loop_clauses = *not_loop_clauses = NULL_TREE;
   for (; clauses ; clauses = next)
@@ -709,27 +708,29 @@ c_oacc_split_loop_clauses (tree clauses, tree *not_loop_clauses)
 
       switch (OMP_CLAUSE_CODE (clauses))
         {
+	  /* Loop clauses.  */
 	case OMP_CLAUSE_COLLAPSE:
-	case OMP_CLAUSE_REDUCTION:
+	case OMP_CLAUSE_TILE:
 	case OMP_CLAUSE_GANG:
 	case OMP_CLAUSE_WORKER:
 	case OMP_CLAUSE_VECTOR:
 	case OMP_CLAUSE_AUTO:
 	case OMP_CLAUSE_SEQ:
+	case OMP_CLAUSE_INDEPENDENT:
+	case OMP_CLAUSE_PRIVATE:
 	  OMP_CLAUSE_CHAIN (clauses) = loop_clauses;
 	  loop_clauses = clauses;
 	  break;
 
-	case OMP_CLAUSE_PRIVATE:
-	  {
-	    tree nc = build_omp_clause (OMP_CLAUSE_LOCATION (clauses),
-					OMP_CLAUSE_CODE (clauses));
-	    OMP_CLAUSE_DECL (nc) = OMP_CLAUSE_DECL (clauses);
-	    OMP_CLAUSE_CHAIN (nc) = loop_clauses;
-	    loop_clauses = nc;
-	  }
-	  /* FALLTHRU */
+	  /* Reductions belong in both constructs.  */
+	case OMP_CLAUSE_REDUCTION:
+	  t = copy_node (clauses);
+	  OMP_CLAUSE_CHAIN (t) = loop_clauses;
+	  loop_clauses = t;
+
+	  /* FIXME: device_type */
 
+	  /* Parallel/kernels clauses.  */
 	default:
 	  OMP_CLAUSE_CHAIN (clauses) = *not_loop_clauses;
 	  *not_loop_clauses = clauses;
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index fa70055..96c1bdc 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -10627,11 +10627,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,
-			     bool only_none = false)
+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;
@@ -10652,7 +10654,7 @@ c_parser_omp_clause_default (c_parser *parser, tree list,
 	  break;
 
 	case 's':
-	  if (strcmp ("shared", p) != 0 || only_none)
+	  if (strcmp ("shared", p) != 0 || is_oacc)
 	    goto invalid_kind;
 	  kind = OMP_CLAUSE_DEFAULT_SHARED;
 	  break;
@@ -10666,7 +10668,7 @@ c_parser_omp_clause_default (c_parser *parser, tree list,
   else
     {
     invalid_kind:
-      if (only_none)
+      if (is_oacc)
 	c_parser_error (parser, "expected %<none%>");
       else
 	c_parser_error (parser, "expected %<none%> or %<shared%>");
@@ -11326,7 +11328,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;
@@ -11571,66 +11576,58 @@ c_parser_oacc_clause_device_type (c_parser *parser, omp_clause_mask mask,
 static tree
 c_parser_oacc_clause_tile (c_parser *parser, tree list)
 {
-  tree c, num = error_mark_node;
-  HOST_WIDE_INT n;
-  location_t loc;
+  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;
-    }
+    return list;
 
   do
     {
-      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);
-	  num = integer_minus_one_node;
+	  expr = integer_minus_one_node;
 	}
       else
 	{
-	  num = c_parser_expr_no_commas (parser, NULL).value;
+	  expr_loc = c_parser_peek_token (parser)->location;
+	  expr = c_parser_expr_no_commas (parser, NULL).value;
 
-	  if (num == error_mark_node)
+	  if (expr == error_mark_node)
 	    {
 	      c_parser_skip_until_found (parser, CPP_CLOSE_PAREN,
 					 "expected %<)%>");
-	      release_tree_vector (tvec);
 	      return list;
 	    }
 
-	  mark_exp_read (num);
-	  num = c_fully_fold (num, false, NULL);
-
-	  if (!INTEGRAL_TYPE_P (TREE_TYPE (num))
-	      || !tree_fits_shwi_p (num)
-	      || (n = tree_to_shwi (num)) <= 0
-	      || (int) n != n)
+	  if (!INTEGRAL_TYPE_P (TREE_TYPE (expr)))
 	    {
-	      error_at (loc,
-			"tile argument needs positive constant integer "
-			"expression");
-	      release_tree_vector (tvec);
-	      c_parser_skip_until_found (parser, CPP_CLOSE_PAREN,
-					 "expected %<)%>");
+	      c_parser_error (parser, "%<tile%> value must be integral");
 	      return list;
 	    }
-	}
 
-      if (num == error_mark_node)
-	{
-	  error_at (loc, "expected positive integer or %<)%>");
-	  release_tree_vector (tvec);
-	  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, num);
+      tile = tree_cons (NULL_TREE, expr, tile);
       if (c_parser_next_token_is (parser, CPP_COMMA))
 	c_parser_consume_token (parser);
     }
@@ -11640,10 +11637,9 @@ c_parser_oacc_clause_tile (c_parser *parser, tree list)
   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);
   return c;
 }
 
@@ -12965,7 +12961,7 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
 	  break;
 	case PRAGMA_OACC_CLAUSE_SEQ:
 	  clauses = c_parser_oacc_simple_clause (parser, OMP_CLAUSE_SEQ,
-						 clauses);
+						clauses);
 	  c_name = "seq";
 	  break;
 	case PRAGMA_OACC_CLAUSE_TILE:
@@ -13060,7 +13056,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:
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index f194b98..fb9b4d0 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -29727,7 +29727,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;
@@ -29864,69 +29867,33 @@ cp_parser_oacc_clause_device_type (cp_parser *parser, omp_clause_mask mask,
    tile ( size-expr-list ) */
 
 static tree
-cp_parser_oacc_clause_tile (cp_parser *parser, tree list, location_t here)
+cp_parser_oacc_clause_tile (cp_parser *parser, location_t clause_loc, tree list)
 {
-  tree c, num = error_mark_node;
-  HOST_WIDE_INT n;
-  location_t loc;
+  tree c, expr = error_mark_node;
   tree tile = NULL_TREE;
-  vec<tree, va_gc> *tvec = make_tree_vector ();
 
-  check_no_duplicate_clause (list, OMP_CLAUSE_TILE, "tile", here);
+  check_no_duplicate_clause (list, OMP_CLAUSE_TILE, "tile", clause_loc);
 
-  loc = cp_lexer_peek_token (parser->lexer)->location;
   if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
-    {
-      release_tree_vector (tvec);
-      return list;
-    }
+    return list;
 
   do
     {
-      if (cp_lexer_next_token_is (parser->lexer, CPP_MULT))
+      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);
-	  num = integer_minus_one_node;
+	  expr = integer_minus_one_node;
 	}
       else
-	{
-	  bool non_constant = false;
-	  num = cp_parser_constant_expression (parser, true, &non_constant);
+	expr = cp_parser_assignment_expression (parser, NULL, false, false);
 
-	  if (num == error_mark_node)
-	    {
-	      cp_parser_skip_to_closing_parenthesis (parser, true, false,
-						     true);
-	      release_tree_vector (tvec);
-	      return list;
-	    }
-
-	  num = fold_non_dependent_expr (num);
-
-	  if (non_constant
-	      || !INTEGRAL_TYPE_P (TREE_TYPE (num))
-	      || !tree_fits_shwi_p (num)
-	      || (n = tree_to_shwi (num)) <= 0
-	      || (int) n != n)
-	    {
-	      error_at (loc,
-			"tile argument needs positive constant integer "
-			"expression");
-	      release_tree_vector (tvec);
-	      cp_parser_skip_to_closing_parenthesis (parser, true, false,
-						     true);
-	      return list;
-	    }
-	}
+      if (expr == error_mark_node)
+	return list;
 
-      if (num == error_mark_node)
-	{
-	  error_at (loc, "expected positive integer or %<)%>");
-	  release_tree_vector (tvec);
-	  return list;
-	}
+      tile = tree_cons (NULL_TREE, expr, tile);
 
-      vec_safe_push (tvec, num);
       if (cp_lexer_next_token_is (parser->lexer, CPP_COMMA))
 	cp_lexer_consume_token (parser->lexer);
     }
@@ -29935,11 +29902,10 @@ cp_parser_oacc_clause_tile (cp_parser *parser, tree list, location_t here)
   /* Consume the trailing ')'.  */
   cp_lexer_consume_token (parser->lexer);
 
-  c = build_omp_clause (loc, OMP_CLAUSE_TILE);
-  tile = build_tree_list_vec (tvec);
+  c = build_omp_clause (clause_loc, OMP_CLAUSE_TILE);
+  tile = nreverse (tile);
   OMP_CLAUSE_TILE_LIST (c) = tile;
   OMP_CLAUSE_CHAIN (c) = list;
-  release_tree_vector (tvec);
   return c;
 }
 
@@ -30051,11 +30017,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, bool is_omp)
+			      location_t location, bool is_oacc)
 {
   enum omp_clause_default_kind kind = OMP_CLAUSE_DEFAULT_UNSPECIFIED;
   tree c;
@@ -30076,7 +30045,7 @@ cp_parser_omp_clause_default (cp_parser *parser, tree list,
 	  break;
 
 	case 's':
-	  if (strcmp ("shared", p) != 0 || !is_omp)
+	  if (strcmp ("shared", p) != 0 || is_oacc)
 	    goto invalid_kind;
 	  kind = OMP_CLAUSE_DEFAULT_SHARED;
 	  break;
@@ -30090,10 +30059,10 @@ cp_parser_omp_clause_default (cp_parser *parser, tree list,
   else
     {
     invalid_kind:
-      if (is_omp)
-	cp_parser_error (parser, "expected %<none%> or %<shared%>");
-      else
+      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))
@@ -31656,8 +31625,7 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
 	  c_name = "delete";
 	  break;
 	case PRAGMA_OMP_CLAUSE_DEFAULT:
-	  clauses = cp_parser_omp_clause_default (parser, clauses, here,
-						  false);
+	  clauses = cp_parser_omp_clause_default (parser, clauses, here, true);
 	  c_name = "default";
 	  break;
 	case PRAGMA_OACC_CLAUSE_DEVICE:
@@ -31758,7 +31726,7 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
 	  c_name = "seq";
 	  break;
 	case PRAGMA_OACC_CLAUSE_TILE:
-	  clauses = cp_parser_oacc_clause_tile (parser, clauses, here);
+	  clauses = cp_parser_oacc_clause_tile (parser, here, clauses);
 	  c_name = "tile";
 	  break;
 	case PRAGMA_OACC_CLAUSE_USE_DEVICE:
@@ -31858,7 +31826,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, true);
+						  token->location, false);
 	  c_name = "default";
 	  break;
 	case PRAGMA_OMP_CLAUSE_FINAL:
@@ -35083,7 +35051,6 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_TYPE)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FIRSTPRIVATE)       	\
-	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_GANG)                \
 	| (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)		\
diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c
index 36951f0..1a45d8d 100644
--- a/gcc/cp/pt.c
+++ b/gcc/cp/pt.c
@@ -14398,7 +14398,6 @@ tsubst_omp_clauses (tree clauses, bool declare_simd, bool allow_fields,
 	case OMP_CLAUSE_NUM_GANGS:
 	case OMP_CLAUSE_NUM_WORKERS:
 	case OMP_CLAUSE_VECTOR_LENGTH:
-	case OMP_CLAUSE_GANG:
 	case OMP_CLAUSE_WORKER:
 	case OMP_CLAUSE_VECTOR:
 	case OMP_CLAUSE_ASYNC:
@@ -14427,7 +14426,7 @@ tsubst_omp_clauses (tree clauses, bool declare_simd, bool allow_fields,
 	    = tsubst_omp_clause_decl (OMP_CLAUSE_DECL (oc), args, complain,
 				      in_decl);
 	  break;
-	case OMP_CLAUSE_LINEAR:
+	case OMP_CLAUSE_GANG:
 	case OMP_CLAUSE_ALIGNED:
 	  OMP_CLAUSE_DECL (nc)
 	    = tsubst_omp_clause_decl (OMP_CLAUSE_DECL (oc), args, complain,
@@ -14460,9 +14459,21 @@ tsubst_omp_clauses (tree clauses, bool declare_simd, bool allow_fields,
 	case OMP_CLAUSE_INDEPENDENT:
 	case OMP_CLAUSE_AUTO:
 	case OMP_CLAUSE_SEQ:
-	case OMP_CLAUSE_TILE:
 	case OMP_CLAUSE_DEVICE_TYPE:
 	  break;
+	case OMP_CLAUSE_TILE:
+	  {
+	    tree lnc, loc;
+	    for (lnc = OMP_CLAUSE_TILE_LIST (nc),
+		   loc = OMP_CLAUSE_TILE_LIST (oc);
+		 loc;
+		 loc = TREE_CHAIN (loc), lnc = TREE_CHAIN (lnc))
+	      {
+		TREE_VALUE (lnc) = tsubst_expr (TREE_VALUE (loc), args,
+						complain, in_decl, false);
+	      }
+	  }
+	  break;
 	default:
 	  gcc_unreachable ();
 	}
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index edcc2f4..e87a906 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -6790,7 +6790,43 @@ finish_omp_clauses (tree clauses, bool is_oacc, bool allow_fields,
 	case OMP_CLAUSE_SEQ:
 	case OMP_CLAUSE_BIND:
 	case OMP_CLAUSE_NOHOST:
+	  break;
+
 	case OMP_CLAUSE_TILE:
+	  for (tree list = OMP_CLAUSE_TILE_LIST (c); !remove && list;
+	       list = TREE_CHAIN (list))
+	    {
+	      t = TREE_VALUE (list);
+
+	      if (t == error_mark_node)
+		remove = true;
+	      else if (!type_dependent_expression_p (t)
+		       && !INTEGRAL_TYPE_P (TREE_TYPE (t)))
+		{
+		  error ("%<tile%> value must be integral");
+		  remove = true;
+		}
+	      else
+		{
+		  t = mark_rvalue_use (t);
+		  if (!processing_template_decl)
+		    {
+		      t = maybe_constant_value (t);
+		      if (TREE_CODE (t) == INTEGER_CST
+			  && tree_int_cst_sgn (t) != 1
+			  && t != integer_minus_one_node)
+			{
+			  warning_at (OMP_CLAUSE_LOCATION (c), 0,
+				      "%<tile%> value must be positive");
+			  t = integer_one_node;
+			}
+		    }
+		  t = fold_build_cleanup_point_expr (TREE_TYPE (t), t);
+		}
+
+		/* Update list item.  */
+	      TREE_VALUE (list) = t;
+	    }
 	  break;
 
 	case OMP_CLAUSE_DEVICE_TYPE:
diff --git a/gcc/testsuite/c-c++-common/goacc/combined-directives.c b/gcc/testsuite/c-c++-common/goacc/combined-directives.c
new file mode 100644
index 0000000..c387285
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/combined-directives.c
@@ -0,0 +1,119 @@
+// { dg-do compile }
+// { dg-options "-fopenacc -fdump-tree-gimple" }
+
+// This error is temporary.  Remove when support is added for these clauses
+// in the middle end.  Also remove the comments from the reduction test
+// after the FE learns that reduction variables may appear in data clauses too.
+// { dg-prune-output "sorry, unimplemented" }
+
+void
+test ()
+{
+  int a[100], i, j, z;
+
+  // acc parallel
+
+  #pragma acc parallel loop collapse (2)
+  for (i = 0; i < 100; i++)
+    for (j = 0; j < 10; j++)
+      ;
+
+  #pragma acc parallel loop gang
+  for (i = 0; i < 100; i++)
+    ;
+
+  #pragma acc parallel loop worker
+  for (i = 0; i < 100; i++)
+    for (j = 0; j < 10; j++)
+      ;
+
+  #pragma acc parallel loop vector
+  for (i = 0; i < 100; i++)
+    for (j = 0; j < 10; j++)
+      ;
+
+  #pragma acc parallel loop seq
+  for (i = 0; i < 100; i++)
+    for (j = 0; j < 10; j++)
+      ;
+
+  #pragma acc parallel loop auto
+  for (i = 0; i < 100; i++)
+    for (j = 0; j < 10; j++)
+      ;
+
+  #pragma acc parallel loop tile (2, 3)
+  for (i = 0; i < 100; i++)
+    for (j = 0; j < 10; j++)
+      ;
+
+  #pragma acc parallel loop independent
+  for (i = 0; i < 100; i++)
+    ;
+
+  #pragma acc parallel loop private (z)
+  for (i = 0; i < 100; i++)
+    z = 0;
+
+//  #pragma acc parallel loop reduction (+:z) copy (z)
+//  for (i = 0; i < 100; i++)
+//    ;
+
+  // acc kernels
+
+  #pragma acc kernels loop collapse (2)
+  for (i = 0; i < 100; i++)
+    for (j = 0; j < 10; j++)
+      ;
+
+  #pragma acc kernels loop gang
+  for (i = 0; i < 100; i++)
+    ;
+
+  #pragma acc kernels loop worker
+  for (i = 0; i < 100; i++)
+    for (j = 0; j < 10; j++)
+      ;
+
+  #pragma acc kernels loop vector
+  for (i = 0; i < 100; i++)
+    for (j = 0; j < 10; j++)
+      ;
+
+  #pragma acc kernels loop seq
+  for (i = 0; i < 100; i++)
+    for (j = 0; j < 10; j++)
+      ;
+
+  #pragma acc kernels loop auto
+  for (i = 0; i < 100; i++)
+    for (j = 0; j < 10; j++)
+      ;
+
+  #pragma acc kernels loop tile (2, 3)
+  for (i = 0; i < 100; i++)
+    for (j = 0; j < 10; j++)
+      ;
+
+  #pragma acc kernels loop independent
+  for (i = 0; i < 100; i++)
+    ;
+
+  #pragma acc kernels loop private (z)
+  for (i = 0; i < 100; i++)
+    z = 0;
+
+//  #pragma acc kernels loop reduction (+:z) copy (z)
+//  for (i = 0; i < 100; i++)
+//    ;
+}
+
+// { dg-final { scan-tree-dump-times "acc loop collapse.2. private.j. private.i" 2 "gimple" } }
+// { dg-final { scan-tree-dump-times "acc loop gang" 2 "gimple" } }
+// { dg-final { scan-tree-dump-times "acc loop worker" 2 "gimple" } }
+// { dg-final { scan-tree-dump-times "acc loop vector" 2 "gimple" } }
+// { dg-final { scan-tree-dump-times "acc loop seq" 2 "gimple" } }
+// { dg-final { scan-tree-dump-times "acc loop auto" 2 "gimple" } }
+// { dg-final { scan-tree-dump-times "acc loop tile.2, 3" 2 "gimple" } }
+// { dg-final { scan-tree-dump-times "acc loop independent private.i" 2 "gimple" } }
+// { dg-final { scan-tree-dump-times "private.z" 2 "gimple" } }
diff --git a/gcc/testsuite/c-c++-common/goacc/loop-clauses.c b/gcc/testsuite/c-c++-common/goacc/loop-clauses.c
new file mode 100644
index 0000000..97b8786
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/loop-clauses.c
@@ -0,0 +1,86 @@
+/* { dg-do compile } */
+
+/* { dg-prune-output "sorry, unimplemented" } */
+
+int
+main ()
+{
+  int i, j;
+
+#pragma acc parallel firstprivate (j) private (i)
+  {
+#pragma acc loop seq
+    for (i = 0; i < 10; i++)
+      { }
+  }
+
+#pragma acc parallel default (none)
+  {
+#pragma acc loop auto private (j)
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop gang
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop gang(static:5)
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop gang(static:*)
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop vector
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop worker
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop auto
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop independent
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop seq
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop gang worker vector
+    for (i = 0; i < 10; i++)
+      { }
+  }
+
+#pragma acc kernels default (none)
+  {
+#pragma acc loop auto
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop gang (num:5)
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop gang(static:5)
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop gang(static:*)
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop vector(length:10)
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop worker(num:5)
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop auto
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop independent
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop seq
+    for (i = 0; i < 10; i++)
+      { }
+#pragma acc loop gang worker vector
+    for (i = 0; i < 10; i++)
+      { }
+  }
+
+  return 0;
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/loop-shape.c b/gcc/testsuite/c-c++-common/goacc/loop-shape.c
index b6d3156..9708f7b 100644
--- a/gcc/testsuite/c-c++-common/goacc/loop-shape.c
+++ b/gcc/testsuite/c-c++-common/goacc/loop-shape.c
@@ -8,6 +8,7 @@ int main ()
   int i;
   int v = 32, w = 19;
   int length = 1, num = 5;
+  int *abc;
 
   /* Valid uses.  */
 
@@ -199,12 +200,12 @@ int main ()
     ;
 
   #pragma acc kernels
-  #pragma acc loop gang(static: * abc) /* { dg-error "expected '.' before" } */
+  #pragma acc loop gang(static: * abc)
   for (i = 0; i < 10; i++)
     ;
 
   #pragma acc kernels
-  #pragma acc loop gang(static:*num:1) /* { dg-error "expected '.' before" } */
+  #pragma acc loop gang(static:*num:1) /* { dg-error "" } */
   for (i = 0; i < 10; i++)
     ;
 
diff --git a/gcc/testsuite/c-c++-common/goacc/loop-tile-k1.c b/gcc/testsuite/c-c++-common/goacc/loop-tile-k1.c
index 9b70193..45cbe2d 100644
--- a/gcc/testsuite/c-c++-common/goacc/loop-tile-k1.c
+++ b/gcc/testsuite/c-c++-common/goacc/loop-tile-k1.c
@@ -35,10 +35,10 @@ kern (void)
 	for (j = 0; j < 10; i++)
 	  { }
       }
-#pragma acc loop tile(-1) // { dg-error "tile argument needs positive constant integer" }
+#pragma acc loop tile(-2) // { dg-warning "'tile' value must be positive" }
     for (i = 0; i < 10; i++)
       { }
-#pragma acc loop tile(i) // { dg-error "tile argument needs positive constant integer" }
+#pragma acc loop tile(i)
     for (i = 0; i < 10; i++)
       { }
 #pragma acc loop tile(2, 2, 1)
@@ -93,10 +93,10 @@ void k3 (void)
       for (j = 1; j < 10; j++)
 	{ }
     }
-#pragma acc kernels loop tile(-1) // { dg-error "tile argument needs positive constant integer expression" }
+#pragma acc kernels loop tile(-2) // { dg-warning "'tile' value must be positive" }
   for (i = 1; i < 10; i++)
     { }
-#pragma acc kernels loop tile(i) // { dg-error "tile argument needs positive constant integer expression" }
+#pragma acc kernels loop tile(i)
   for (i = 1; i < 10; i++)
     { }
 #pragma acc kernels loop tile(2, 2, 1)
diff --git a/gcc/testsuite/c-c++-common/goacc/loop-tile-p1.c b/gcc/testsuite/c-c++-common/goacc/loop-tile-p1.c
index 41c0b24..665bc15 100644
--- a/gcc/testsuite/c-c++-common/goacc/loop-tile-p1.c
+++ b/gcc/testsuite/c-c++-common/goacc/loop-tile-p1.c
@@ -26,10 +26,10 @@ void par (void)
 	for (j = 1; j < 10; j++)
 	  { }
       }
-#pragma acc loop tile(-1) // { dg-error "tile argument needs positive constant integer expression" }
+#pragma acc loop tile(-2) // { dg-warning "'tile' value must be positive" }
     for (i = 1; i < 10; i++)
       { }
-#pragma acc loop tile(i) // { dg-error "tile argument needs positive constant integer expression" }
+#pragma acc loop tile(i)
     for (i = 1; i < 10; i++)
       { }
 #pragma acc loop tile(2, 2, 1)
@@ -87,10 +87,10 @@ void p3 (void)
       for (j = 1; j < 10; j++)
 	{ }
     }
-#pragma acc parallel loop tile(-1) // { dg-error "tile argument needs positive constant integer expression" }
+#pragma acc parallel loop tile(-2) // { dg-warning "'tile' value must be positive" }
   for (i = 1; i < 10; i++)
     { }
-#pragma acc parallel loop tile(i) // { dg-error "tile argument needs positive constant integer expression" }
+#pragma acc parallel loop tile(i)
   for (i = 1; i < 10; i++)
     { }
 #pragma acc parallel loop tile(2, 2, 1)
diff --git a/gcc/testsuite/c-c++-common/goacc/tile.c b/gcc/testsuite/c-c++-common/goacc/tile.c
index 57de2a8..2a81427 100644
--- a/gcc/testsuite/c-c++-common/goacc/tile.c
+++ b/gcc/testsuite/c-c++-common/goacc/tile.c
@@ -1,7 +1,9 @@
+/* { dg-do compile } */
+
 int
 main ()
 {
-  int i;
+  int i, *a, b;
 
 #pragma acc parallel loop tile (10)
   for (i = 0; i < 100; i++)
@@ -15,7 +17,55 @@ main ()
   for (i = 0; i < 100; i++)
     ;
 
-#pragma acc parallel loop tile (10, *, i) /* { dg-error "positive constant integer expression" } */
+#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++)
     ;