In working on implementing tile, I discovered the parsing of tile clauses was at
best confused. (The spec didn't help by being somewhat vague and not
considering interaction with the collapse clause)
Fixed with the attached patch.
nathan
2016-09-23 Nathan Sidwell <nat...@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
+}
