[gomp4] Fix tile parsing

2016-09-23  Nathan Sidwell  <nathan@codesourcery.com>

	c/
	* c-parser.c (c_parser_omp_clause_collapse): Disallow tile.
	(c_parser_oacc_clause_tile): Disallow collapse. Fix parsing and
	semantic checking.

	cp/
	* parser.c (cp_parser_oacc_clause_tile): Disallow collapse.  Fix
	parsing.  Parse constant expression. Remove semantic checking.
	(cp_parser_omp_clause_collapse): Disallow tile.
	* pt.c (tsubst_omp_clauses): Require integral constant expression
	for COLLAPSE and TILE.  Remove broken TILE subst.
	* semantics.c (finish_omp_clauses): Correct TILE semantic check.

	testsuite/
	* c-c++-common/goacc/tile.c: Include stdbool, fix expected errors.
	* g++.dg/goacc/template.C: Test tile subst.  Adjust erroneous
	uses.
	* g++.dg/goacc/tile-1.C: Check tile subst.

Index: c/c-parser.c
===================================================================
--- c/c-parser.c	(revision 240420)
+++ c/c-parser.c	(working copy)
@@ -10882,6 +10882,7 @@  c_parser_omp_clause_collapse (c_parser *
   location_t loc;
 
   check_no_duplicate_clause (list, OMP_CLAUSE_COLLAPSE, "collapse");
+  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 %<(%>"))
@@ -11905,10 +11906,11 @@  static tree
 c_parser_oacc_clause_tile (c_parser *parser, tree list)
 {
   tree c, expr = error_mark_node;
-  location_t loc, expr_loc;
+  location_t loc;
   tree tile = NULL_TREE;
 
   check_no_duplicate_clause (list, OMP_CLAUSE_TILE, "tile");
+  check_no_duplicate_clause (list, OMP_CLAUSE_COLLAPSE, "collapse");
 
   loc = c_parser_peek_token (parser)->location;
   if (!c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>"))
@@ -11916,16 +11918,19 @@  c_parser_oacc_clause_tile (c_parser *par
 
   do
     {
+      if (tile && !c_parser_require (parser, CPP_COMMA, "expected %<,%>"))
+	return list;
+
       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;
+	  expr = integer_zero_node;
 	}
       else
 	{
-	  expr_loc = c_parser_peek_token (parser)->location;
+	  location_t expr_loc = c_parser_peek_token (parser)->location;
 	  expr = c_parser_expr_no_commas (parser, NULL).value;
 
 	  if (expr == error_mark_node)
@@ -11935,29 +11940,21 @@  c_parser_oacc_clause_tile (c_parser *par
 	      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)
+	  if (!INTEGRAL_TYPE_P (TREE_TYPE (expr))
+	      || TREE_CODE (expr) != INTEGER_CST
+	      || !tree_fits_shwi_p (expr)
+	      || tree_to_shwi (expr) <= 0)
 	    {
-	      warning_at (expr_loc, 0,"%<tile%> value must be positive");
-	      expr = integer_one_node;
+	      error_at (expr_loc, "%<tile%> argument needs positive"
+			" integral constant");
+	      expr = integer_zero_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));
 
Index: cp/parser.c
===================================================================
--- cp/parser.c	(revision 240420)
+++ cp/parser.c	(working copy)
@@ -30434,30 +30434,33 @@  cp_parser_oacc_clause_tile (cp_parser *p
   tree c, expr = error_mark_node;
   tree tile = NULL_TREE;
 
+  /* Collapse and tile are mutually exclusive.  (The spec doesn't say
+     so, but the spec authors never considered such a case and have
+     differing opinions on what it might mean, including 'not
+     allowed'.)  */
   check_no_duplicate_clause (list, OMP_CLAUSE_TILE, "tile", clause_loc);
+  check_no_duplicate_clause (list, OMP_CLAUSE_COLLAPSE, "collapse",
+			     clause_loc);
 
   if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
     return list;
 
   do
     {
+      if (tile && !cp_parser_require (parser, CPP_COMMA, RT_COMMA))
+	return list;
+      
       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;
+	  expr = integer_zero_node;
 	}
       else
-	expr = cp_parser_assignment_expression (parser, NULL, false, false);
-
-      if (expr == error_mark_node)
-	return list;
+	expr = cp_parser_constant_expression (parser);
 
       tile = tree_cons (NULL_TREE, expr, tile);
-
-      if (cp_lexer_next_token_is (parser->lexer, CPP_COMMA))
-	cp_lexer_consume_token (parser->lexer);
     }
   while (cp_lexer_next_token_is_not (parser->lexer, CPP_CLOSE_PAREN));
 
@@ -30571,6 +30574,7 @@  cp_parser_omp_clause_collapse (cp_parser
     }
 
   check_no_duplicate_clause (list, OMP_CLAUSE_COLLAPSE, "collapse", location);
+  check_no_duplicate_clause (list, OMP_CLAUSE_TILE, "tile", location);
   c = build_omp_clause (loc, OMP_CLAUSE_COLLAPSE);
   OMP_CLAUSE_CHAIN (c) = list;
   OMP_CLAUSE_COLLAPSE_EXPR (c) = num;
Index: cp/pt.c
===================================================================
--- cp/pt.c	(revision 240420)
+++ cp/pt.c	(working copy)
@@ -14543,6 +14543,7 @@  tsubst_omp_clauses (tree clauses, enum c
       nc = copy_node (oc);
       OMP_CLAUSE_CHAIN (nc) = new_clauses;
       new_clauses = nc;
+      bool needs_ice = false;
 
       switch (OMP_CLAUSE_CODE (nc))
 	{
@@ -14572,10 +14573,16 @@  tsubst_omp_clauses (tree clauses, enum c
 	    = tsubst_omp_clause_decl (OMP_CLAUSE_DECL (oc), args, complain,
 				      in_decl);
 	  break;
+	case OMP_CLAUSE_COLLAPSE:
+	case OMP_CLAUSE_TILE:
+	  /* These clauses really need a positive integral constant
+	     expression, but we don't have a predicate for that
+	     (yet).  */
+	  needs_ice = true;
+	  /* FALLTHRU */
 	case OMP_CLAUSE_IF:
 	case OMP_CLAUSE_NUM_THREADS:
 	case OMP_CLAUSE_SCHEDULE:
-	case OMP_CLAUSE_COLLAPSE:
 	case OMP_CLAUSE_FINAL:
 	case OMP_CLAUSE_DEVICE:
 	case OMP_CLAUSE_DIST_SCHEDULE:
@@ -14596,8 +14603,8 @@  tsubst_omp_clauses (tree clauses, enum c
 	case OMP_CLAUSE_ASYNC:
 	case OMP_CLAUSE_WAIT:
 	  OMP_CLAUSE_OPERAND (nc, 0)
-	    = tsubst_expr (OMP_CLAUSE_OPERAND (oc, 0), args, complain, 
-			   in_decl, /*integral_constant_expression_p=*/false);
+	    = tsubst_expr (OMP_CLAUSE_OPERAND (oc, 0), args, complain, in_decl,
+			   /*integral_constant_expression_p=*/needs_ice);
 	  break;
 	case OMP_CLAUSE_REDUCTION:
 	  if (OMP_CLAUSE_REDUCTION_PLACEHOLDER (oc))
@@ -14667,19 +14674,6 @@  tsubst_omp_clauses (tree clauses, enum c
 	case OMP_CLAUSE_SEQ:
 	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;
 	case OMP_CLAUSE_BIND:
 	case OMP_CLAUSE_NOHOST:
 	default:
Index: cp/semantics.c
===================================================================
--- cp/semantics.c	(revision 240420)
+++ cp/semantics.c	(working copy)
@@ -7063,7 +7063,8 @@  finish_omp_clauses (tree clauses, enum c
 	      else if (!type_dependent_expression_p (t)
 		       && !INTEGRAL_TYPE_P (TREE_TYPE (t)))
 		{
-		  error ("%<tile%> value must be integral");
+		  error_at (OMP_CLAUSE_LOCATION (c),
+			    "%<tile%> argument needs integral type");
 		  remove = true;
 		}
 	      else
@@ -7071,14 +7072,16 @@  finish_omp_clauses (tree clauses, enum c
 		  t = mark_rvalue_use (t);
 		  if (!processing_template_decl)
 		    {
+		      /* Zero is used to indicate '*', we permit you
+			 to get there via an ICE of value zero.  */
 		      t = maybe_constant_value (t);
-		      if (TREE_CODE (t) == INTEGER_CST
-			  && tree_int_cst_sgn (t) != 1
-			  && t != integer_minus_one_node)
+		      if (TREE_CODE (t) != INTEGER_CST
+			  || !tree_fits_shwi_p (t)
+			  || tree_to_shwi (t) < 0)
 			{
-			  warning_at (OMP_CLAUSE_LOCATION (c), 0,
-				      "%<tile%> value must be positive");
-			  t = integer_one_node;
+			  error_at (OMP_CLAUSE_LOCATION (c),
+				    "%<tile%> argument needs positive integral constant");
+			  remove = true;
 			}
 		    }
 		  t = fold_build_cleanup_point_expr (TREE_TYPE (t), t);
Index: testsuite/c-c++-common/goacc/tile.c
===================================================================
--- testsuite/c-c++-common/goacc/tile.c	(revision 240420)
+++ testsuite/c-c++-common/goacc/tile.c	(working copy)
@@ -1,3 +1,5 @@ 
+#include <stdbool.h>
+
 int
 main ()
 {
@@ -15,7 +17,7 @@  main ()
   for (i = 0; i < 100; i++)
     ;
 
-#pragma acc parallel loop tile (10, *, i)
+#pragma acc parallel loop tile (10, *, i) // { dg-error "" }
   for (i = 0; i < 100; i++)
     ;
 
@@ -35,35 +37,35 @@  main ()
   for (i = 0; i < 100; i++)
     ;
 
-#pragma acc parallel loop tile (1.1) // { dg-error "'tile' value must be integral" }
+#pragma acc parallel loop tile (1.1) // { dg-error "'tile' argument needs" }
   for (i = 0; i < 100; i++)
     ;
 
-#pragma acc parallel loop tile (-3) // { dg-warning "'tile' value must be positive" }
+#pragma acc parallel loop tile (-3) // { dg-error "'tile' argument needs" }
   for (i = 0; i < 100; i++)
     ;
 
-#pragma acc parallel loop tile (10,-3) // { dg-warning "'tile' value must be positive" }
+#pragma acc parallel loop tile (10,-3) // { dg-error "'tile' argument needs" }
   for (i = 0; i < 100; i++)
     ;
 
-#pragma acc parallel loop tile (-100,10,5) // { dg-warning "'tile' value must be positive" }
+#pragma acc parallel loop tile (-100,10,5) // { dg-error "'tile' argument needs" }
   for (i = 0; i < 100; i++)
     ;
 
-#pragma acc parallel loop tile (1,2.0,true) // { dg-error "" }
+#pragma acc parallel loop tile (1,true)
   for (i = 0; i < 100; i++)
     ;
 
-#pragma acc parallel loop tile (*a, 1)
+#pragma acc parallel loop tile (*a, 1) // { dg-error "" }
   for (i = 0; i < 100; i++)
     ;
 
-#pragma acc parallel loop tile (1, *a, b)
+#pragma acc parallel loop tile (1, b) // { dg-error "" }
   for (i = 0; i < 100; i++)
     ;
 
-#pragma acc parallel loop tile (b, 1, *a)
+#pragma acc parallel loop tile (b, 1) // { dg-error "" }
   for (i = 0; i < 100; i++)
     ;
 
@@ -95,10 +97,10 @@  void par (void)
 	for (j = 1; j < 10; j++)
 	  { }
       }
-#pragma acc loop tile(-2) // { dg-warning "'tile' value must be positive" }
+#pragma acc loop tile(-2)  // { dg-error "'tile' argument needs" }
     for (i = 1; i < 10; i++)
       { }
-#pragma acc loop tile(i)
+#pragma acc loop tile(i)  // { dg-error "" }
     for (i = 1; i < 10; i++)
       { }
 #pragma acc loop tile(2, 2, 1)
@@ -156,24 +158,21 @@  void p3 (void)
       for (j = 1; j < 10; j++)
 	{ }
     }
-#pragma acc parallel loop tile(-2) // { dg-warning "'tile' value must be positive" }
+#pragma acc parallel loop tile(-2)   // { dg-error "'tile' argument needs" }
   for (i = 1; i < 10; i++)
     { }
-#pragma acc parallel loop tile(i)
+#pragma acc parallel loop tile(i)   // { dg-error "" }
   for (i = 1; i < 10; i++)
     { }
 #pragma acc parallel loop tile(2, 2, 1)
   for (i = 1; i < 3; i++)
-    {
-      for (j = 4; j < 6; j++)
-        { }
-    }    
+    for (j = 4; j < 6; j++)
+      for (int k = 1 ; k < 2; k++)
+	;
 #pragma acc parallel loop tile(2, 2)
   for (i = 1; i < 5; i+=2)
-    {
-      for (j = i + 1; j < 7; j++)
-        { }
-    }
+    for (j = i + 1; j < 7; j++)
+      { }
 #pragma acc parallel loop vector tile(*) 
   for (i = 0; i < 10; i++)
     { }
@@ -230,10 +229,10 @@  kern (void)
 	for (j = 0; j < 10; i++)
 	  { }
       }
-#pragma acc loop tile(-2) // { dg-warning "'tile' value must be positive" }
+#pragma acc loop tile(-2) // { dg-error "'tile' argument needs" }
     for (i = 0; i < 10; i++)
       { }
-#pragma acc loop tile(i)
+#pragma acc loop tile(i) // { dg-error "" }
     for (i = 0; i < 10; i++)
       { }
 #pragma acc loop tile(2, 2, 1)
@@ -288,18 +287,17 @@  void k3 (void)
       for (j = 1; j < 10; j++)
 	{ }
     }
-#pragma acc kernels loop tile(-2) // { dg-warning "'tile' value must be positive" }
+#pragma acc kernels loop tile(-2) // { dg-error "'tile' argument needs" }
   for (i = 1; i < 10; i++)
     { }
-#pragma acc kernels loop tile(i)
+#pragma acc kernels loop tile(i) // { dg-error "" }
   for (i = 1; i < 10; i++)
     { }
 #pragma acc kernels loop tile(2, 2, 1)
   for (i = 1; i < 3; i++)
-    {
-      for (j = 4; j < 6; j++)
-	{ }
-    }    
+    for (j = 4; j < 6; j++)
+      for (int k = 1; k < 7; k++)
+	;
 #pragma acc kernels loop tile(2, 2)
   for (i = 1; i < 5; i++)
     {
Index: testsuite/g++.dg/goacc/template.C
===================================================================
--- testsuite/g++.dg/goacc/template.C	(revision 240420)
+++ testsuite/g++.dg/goacc/template.C	(working copy)
@@ -5,7 +5,7 @@  accDouble(int val)
   return val * 2;
 }
 
-template<typename T> T
+template<typename T, int I> T
 oacc_parallel_copy (T a)
 {
   T b = 0;
@@ -36,7 +36,7 @@  oacc_parallel_copy (T a)
       for (int j = 0; j < 5; j++)
 	b = a;
 
-#pragma acc loop auto tile (a, 3)
+#pragma acc loop auto tile (I, 3)
     for (int i = 0; i < a; i++)
       for (int j = 0; j < 5; j++)
 	b = a;
@@ -135,7 +135,7 @@  oacc_kernels_copy (T a)
 int
 main ()
 {
-  int b = oacc_parallel_copy<int> (5);
+  int b = oacc_parallel_copy<int, 4> (5);
   int c = oacc_kernels_copy<int> (5);
 
   return b + c;
Index: testsuite/g++.dg/goacc/tile-1.C
===================================================================
--- testsuite/g++.dg/goacc/tile-1.C	(nonexistent)
+++ testsuite/g++.dg/goacc/tile-1.C	(working copy)
@@ -0,0 +1,16 @@ 
+/*  of tile erroneously clobbered the template, resulting
+    in missing errors and other fun.  */
+
+template <int I>
+void Foo ()
+{
+#pragma acc parallel loop tile(I) // { dg-error "" }
+  for (int ix = 0; ix < 10; ix++)
+    ;
+}
+
+int main ()
+{
+  Foo<1> ();  // OK
+  Foo<-1> (); // error
+}