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.
===================================================================
@@ -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));
===================================================================
@@ -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;
===================================================================
@@ -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:
===================================================================
@@ -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);
===================================================================
@@ -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++)
{
===================================================================
@@ -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;
===================================================================
@@ -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
+}