These are the patches for the C/C++ front-ends, along with the
testsuite patches.
Thanks,
Chung-Lin
2016-XX-XX 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.
* c-parser.c (c_parser_omp_for_loop): Accept tiling constructs.
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.
(cp_parser_omp_for_loop): Deal with tile clause. Don't emit a
parse error about missing for after already emitting one.
Use more conventional for idiom for unbounded loop.
* 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.
(finish_omp_for): Deal with tile clause.
gcc/testsuite/
* c-c++-common/goacc/loop-auto-1.c: Adjust and add additional
case.
* c-c++-common/goacc/loop-auto-2.c: New.
* 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.
* gcc.dg/goacc/loop-processing-1.c: Adjust dg-final pattern.
Index: c/c-parser.c
===================================================================
--- c/c-parser.c (revision 241809)
+++ c/c-parser.c (working copy)
@@ -11010,6 +11010,7 @@ c_parser_omp_clause_collapse (c_parser *parser, tr
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 %<(%>"))
@@ -11920,10 +11921,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 %<(%>"))
@@ -11931,16 +11933,19 @@ c_parser_oacc_clause_tile (c_parser *parser, tree
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;
c_expr cexpr = c_parser_expr_no_commas (parser, NULL);
cexpr = convert_lvalue_to_rvalue (expr_loc, cexpr, false, true);
expr = cexpr.value;
@@ -11952,28 +11957,20 @@ c_parser_oacc_clause_tile (c_parser *parser, tree
return list;
}
- if (!INTEGRAL_TYPE_P (TREE_TYPE (expr)))
- {
- c_parser_error (parser, "%<tile%> value must be integral");
- return list;
- }
-
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));
@@ -14899,11 +14896,17 @@ c_parser_omp_for_loop (location_t loc, c_parser *p
bool fail = false, open_brace_parsed = false;
int i, collapse = 1, ordered = 0, count, nbraces = 0;
location_t for_loc;
+ bool tiling = false;
vec<tree, va_gc> *for_block = make_tree_vector ();
for (cl = clauses; cl; cl = OMP_CLAUSE_CHAIN (cl))
if (OMP_CLAUSE_CODE (cl) == OMP_CLAUSE_COLLAPSE)
collapse = tree_to_shwi (OMP_CLAUSE_COLLAPSE_EXPR (cl));
+ else if (OMP_CLAUSE_CODE (cl) == OMP_CLAUSE_TILE)
+ {
+ tiling = true;
+ collapse = list_length (OMP_CLAUSE_TILE_LIST (cl));
+ }
else if (OMP_CLAUSE_CODE (cl) == OMP_CLAUSE_ORDERED
&& OMP_CLAUSE_ORDERED_EXPR (cl))
{
@@ -14933,7 +14936,7 @@ c_parser_omp_for_loop (location_t loc, c_parser *p
pc = &OMP_CLAUSE_CHAIN (*pc);
}
- gcc_assert (collapse >= 1 && ordered >= 0);
+ gcc_assert (tiling || (collapse >= 1 && ordered >= 0));
count = ordered ? ordered : collapse;
declv = make_tree_vec (count);
Index: cp/pt.c
===================================================================
--- cp/pt.c (revision 241809)
+++ cp/pt.c (working copy)
@@ -14713,6 +14713,7 @@ tsubst_omp_clauses (tree clauses, enum c_omp_regio
nc = copy_node (oc);
OMP_CLAUSE_CHAIN (nc) = new_clauses;
new_clauses = nc;
+ bool needs_ice = false;
switch (OMP_CLAUSE_CODE (nc))
{
@@ -14742,10 +14743,16 @@ tsubst_omp_clauses (tree clauses, enum c_omp_regio
= 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:
@@ -14766,8 +14773,8 @@ tsubst_omp_clauses (tree clauses, enum c_omp_regio
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))
@@ -14836,19 +14843,6 @@ tsubst_omp_clauses (tree clauses, enum c_omp_regio
case OMP_CLAUSE_AUTO:
case OMP_CLAUSE_SEQ:
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 ();
}
Index: cp/semantics.c
===================================================================
--- cp/semantics.c (revision 241809)
+++ cp/semantics.c (working copy)
@@ -7086,7 +7086,8 @@ finish_omp_clauses (tree clauses, enum c_omp_regio
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
@@ -7094,14 +7095,16 @@ finish_omp_clauses (tree clauses, enum c_omp_regio
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);
@@ -8000,11 +8003,19 @@ finish_omp_for (location_t locus, enum tree_code c
gcc_assert (TREE_VEC_LENGTH (declv) == TREE_VEC_LENGTH (incrv));
if (TREE_VEC_LENGTH (declv) > 1)
{
- tree c = find_omp_clause (clauses, OMP_CLAUSE_COLLAPSE);
+ tree c;
+
+ c = find_omp_clause (clauses, OMP_CLAUSE_TILE);
if (c)
- collapse = tree_to_shwi (OMP_CLAUSE_COLLAPSE_EXPR (c));
- if (collapse != TREE_VEC_LENGTH (declv))
- ordered = TREE_VEC_LENGTH (declv);
+ collapse = list_length (OMP_CLAUSE_TILE_LIST (c));
+ else
+ {
+ c = find_omp_clause (clauses, OMP_CLAUSE_COLLAPSE);
+ if (c)
+ collapse = tree_to_shwi (OMP_CLAUSE_COLLAPSE_EXPR (c));
+ if (collapse != TREE_VEC_LENGTH (declv))
+ ordered = TREE_VEC_LENGTH (declv);
+ }
}
for (i = 0; i < TREE_VEC_LENGTH (declv); i++)
{
Index: cp/parser.c
===================================================================
--- cp/parser.c (revision 241809)
+++ cp/parser.c (working copy)
@@ -30885,30 +30885,33 @@ cp_parser_oacc_clause_tile (cp_parser *parser, loc
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);
+ expr = cp_parser_constant_expression (parser);
- if (expr == error_mark_node)
- return list;
-
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));
@@ -31021,6 +31024,7 @@ cp_parser_omp_clause_collapse (cp_parser *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;
@@ -34027,10 +34031,16 @@ cp_parser_omp_for_loop (cp_parser *parser, enum tr
int i, collapse = 1, ordered = 0, count, nbraces = 0;
vec<tree, va_gc> *for_block = make_tree_vector ();
auto_vec<tree, 4> orig_inits;
+ bool tiling = false;
for (cl = clauses; cl; cl = OMP_CLAUSE_CHAIN (cl))
if (OMP_CLAUSE_CODE (cl) == OMP_CLAUSE_COLLAPSE)
collapse = tree_to_shwi (OMP_CLAUSE_COLLAPSE_EXPR (cl));
+ else if (OMP_CLAUSE_CODE (cl) == OMP_CLAUSE_TILE)
+ {
+ tiling = true;
+ collapse = list_length (OMP_CLAUSE_TILE_LIST (cl));
+ }
else if (OMP_CLAUSE_CODE (cl) == OMP_CLAUSE_ORDERED
&& OMP_CLAUSE_ORDERED_EXPR (cl))
{
@@ -34060,7 +34070,7 @@ cp_parser_omp_for_loop (cp_parser *parser, enum tr
pc = &OMP_CLAUSE_CHAIN (*pc);
}
- gcc_assert (collapse >= 1 && ordered >= 0);
+ gcc_assert (tiling || (collapse >= 1 && ordered >= 0));
count = ordered ? ordered : collapse;
declv = make_tree_vec (count);
@@ -34079,13 +34089,15 @@ cp_parser_omp_for_loop (cp_parser *parser, enum tr
if (code != CILK_FOR
&& !cp_lexer_next_token_is_keyword (parser->lexer, RID_FOR))
{
- cp_parser_error (parser, "for statement expected");
+ if (!collapse_err)
+ cp_parser_error (parser, "for statement expected");
return NULL;
}
if (code == CILK_FOR
&& !cp_lexer_next_token_is_keyword (parser->lexer, RID_CILK_FOR))
{
- cp_parser_error (parser, "_Cilk_for statement expected");
+ if (!collapse_err)
+ cp_parser_error (parser, "_Cilk_for statement expected");
return NULL;
}
loc = cp_lexer_consume_token (parser->lexer)->location;
@@ -34245,7 +34257,7 @@ cp_parser_omp_for_loop (cp_parser *parser, enum tr
nested. Hopefully the final version clarifies this.
For now handle (multiple) {'s and empty statements. */
cp_parser_parse_tentatively (parser);
- do
+ for (;;)
{
if (cp_lexer_next_token_is_keyword (parser->lexer, RID_FOR))
break;
@@ -34260,14 +34272,13 @@ cp_parser_omp_for_loop (cp_parser *parser, enum tr
else
{
loc = cp_lexer_peek_token (parser->lexer)->location;
- error_at (loc, "not enough collapsed for loops");
+ error_at (loc, "not enough for loops to collapse");
collapse_err = true;
cp_parser_abort_tentative_parse (parser);
declv = NULL_TREE;
break;
}
}
- while (1);
if (declv)
{
Index: testsuite/c-c++-common/goacc/loop-auto-2.c
===================================================================
--- testsuite/c-c++-common/goacc/loop-auto-2.c (revision 0)
+++ testsuite/c-c++-common/goacc/loop-auto-2.c (revision 0)
@@ -0,0 +1,107 @@
+
+// Tile parititioning
+
+void Ok ()
+{
+#pragma acc parallel num_gangs (10) num_workers(32) vector_length(32)
+ {
+
+#pragma acc loop tile(*) gang vector
+ for (int ix = 0; ix < 10; ix++)
+ {
+ }
+
+#pragma acc loop tile(*)
+ for (int ix = 0; ix < 10; ix++)
+ {
+ }
+
+#pragma acc loop tile(*) gang
+ for (int ix = 0; ix < 10; ix++)
+ {
+ #pragma acc loop vector
+ for (int jx = 0; jx < 10; jx++)
+ ;
+ }
+
+#pragma acc loop tile(*)
+ for (int ix = 0; ix < 10; ix++)
+ {
+ #pragma acc loop vector
+ for (int jx = 0; jx < 10; jx++)
+ ;
+ }
+
+#pragma acc loop gang
+ for (int jx = 0; jx < 10; jx++)
+ {
+#pragma acc loop tile(*) vector
+ for (int ix = 0; ix < 10; ix++)
+ {
+ }
+
+#pragma acc loop tile(*)
+ for (int ix = 0; ix < 10; ix++)
+ {
+ }
+ }
+
+#pragma acc loop tile(*) worker
+ for (int ix = 0; ix < 10; ix++)
+ {
+ #pragma acc loop vector
+ for (int jx = 0; jx < 10; jx++)
+ ;
+ }
+ }
+}
+
+void Bad ()
+{
+#pragma acc parallel num_gangs (10) num_workers(32) vector_length(32)
+ {
+
+#pragma acc loop tile(*) gang vector /* { dg-message "containing loop" } */
+ for (int ix = 0; ix < 10; ix++)
+ {
+#pragma acc loop vector /* { dg-error "uses same" } */
+ for (int jx = 0; jx < 10; jx++)
+ ;
+ }
+
+#pragma acc loop tile(*) gang vector
+ for (int ix = 0; ix < 10; ix++)
+ {
+ #pragma acc loop auto /* { dg-warning "insufficient partitioning" } */
+ for (int jx = 0; jx < 10; jx++)
+ ;
+ }
+
+#pragma acc loop tile(*) auto /* { dg-warning "insufficient partitioning" } */
+ for (int ix = 0; ix < 10; ix++)
+ {
+ #pragma acc loop worker
+ for (int jx = 0; jx < 10; jx++)
+ ;
+ }
+
+#pragma acc loop worker /* { dg-message "containing loop" } */
+ for (int jx = 0; jx < 10; jx++)
+ {
+#pragma acc loop tile(*) gang vector /* { dg-error "incorrectly nested" } */
+ for (int ix = 0; ix < 10; ix++)
+ {
+ }
+
+#pragma acc loop tile(*) vector /* { dg-warning "insufficient partitioning" } */
+ for (int ix = 0; ix < 10; ix++)
+ {
+ }
+
+#pragma acc loop tile(*) /* { dg-warning "insufficient partitioning" } */
+ for (int ix = 0; ix < 10; ix++)
+ {
+ }
+ }
+ }
+}
Index: testsuite/c-c++-common/goacc/tile-2.c
===================================================================
--- testsuite/c-c++-common/goacc/tile-2.c (revision 0)
+++ testsuite/c-c++-common/goacc/tile-2.c (revision 0)
@@ -0,0 +1,21 @@
+int main ()
+{
+#pragma acc parallel
+ {
+#pragma acc loop tile (*,*)
+ for (int ix = 0; ix < 30; ix++)
+ ; /* { dg-error "not enough" } */
+
+#pragma acc loop tile (*,*)
+ for (int ix = 0; ix < 30; ix++)
+ for (int jx = 0; jx < ix; jx++) /* { dg-error "condition expression" } */
+ ;
+
+#pragma acc loop tile (*)
+ for (int ix = 0; ix < 30; ix++)
+ for (int jx = 0; jx < ix; jx++) /* OK */
+ ;
+
+ }
+ return 0;
+}
Index: testsuite/c-c++-common/goacc/loop-auto-1.c
===================================================================
--- testsuite/c-c++-common/goacc/loop-auto-1.c (revision 241809)
+++ testsuite/c-c++-common/goacc/loop-auto-1.c (working copy)
@@ -74,6 +74,21 @@ void Foo ()
for (int kx = 0; kx < 10; kx++) {}
}
}
+
+#pragma acc loop auto
+ for (int ix = 0; ix < 10; ix++)
+ {
+#pragma acc loop auto
+ for (int jx = 0; jx < 10; jx++)
+ {
+#pragma acc loop auto /* { dg-warning "insufficient partitioning" } */
+ for (int kx = 0; kx < 10; kx++)
+ {
+#pragma acc loop auto
+ for (int lx = 0; lx < 10; lx++) {}
+ }
+ }
+ }
}
}
@@ -214,10 +229,10 @@ void Vector (void)
#pragma acc loop auto
for (int ix = 0; ix < 10; ix++) {}
-#pragma acc loop auto
+#pragma acc loop auto /* { dg-warning "insufficient partitioning" } */
for (int ix = 0; ix < 10; ix++)
{
-#pragma acc loop auto /* { dg-warning "insufficient partitioning" } */
+#pragma acc loop auto
for (int jx = 0; jx < 10; jx++) {}
}
}
Index: testsuite/c-c++-common/goacc/tile.c
===================================================================
--- testsuite/c-c++-common/goacc/tile.c (revision 241809)
+++ testsuite/c-c++-common/goacc/tile.c (working copy)
@@ -1,7 +1,9 @@
+#include <stdbool.h>
+
int
main ()
{
- int i, *a, b;
+ int i, j, k, *a, b;
#pragma acc parallel loop tile (10)
for (i = 0; i < 100; i++)
@@ -13,11 +15,14 @@ main ()
#pragma acc parallel loop tile (10, *)
for (i = 0; i < 100; i++)
- ;
+ for (j = 0; j < 100; j++)
+ ;
-#pragma acc parallel loop tile (10, *, i)
+#pragma acc parallel loop tile (10, *, i) // { dg-error "" }
for (i = 0; i < 100; i++)
- ;
+ for (j = 0; j < 100; j++)
+ for (k = 0; k < 100; k++)
+ ;
#pragma acc parallel loop tile // { dg-error "expected '\\\('" }
for (i = 0; i < 100; i++)
@@ -35,37 +40,44 @@ 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++)
- ;
+ for (j = 0; j < 100; j++)
+ ;
-#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++)
- ;
+ for (j = 0; j < 100; j++)
+ for (k = 0; k < 100; k++)
+ ;
-#pragma acc parallel loop tile (1,2.0,true) // { dg-error "" }
+#pragma acc parallel loop tile (1,true)
for (i = 0; i < 100; i++)
- ;
+ for (j = 0; j < 100; j++)
+ ;
-#pragma acc parallel loop tile (*a, 1)
+#pragma acc parallel loop tile (*a, 1) // { dg-error "" }
for (i = 0; i < 100; i++)
- ;
+ for (j = 0; j < 100; j++)
+ ;
-#pragma acc parallel loop tile (1, *a, b)
+#pragma acc parallel loop tile (1, b) // { dg-error "" }
for (i = 0; i < 100; i++)
- ;
+ for (j = 0; j < 100; j++)
+ ;
-#pragma acc parallel loop tile (b, 1, *a)
+#pragma acc parallel loop tile (b, 1) // { dg-error "" }
for (i = 0; i < 100; i++)
- ;
+ for (j = 0; j < 100; j++)
+ ;
return 0;
}
@@ -73,7 +85,7 @@ main ()
void par (void)
{
- int i, j;
+ int i, j, k;
#pragma acc parallel
{
@@ -95,22 +107,22 @@ 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)
for (i = 1; i < 3; i++)
{
for (j = 4; j < 6; j++)
- { }
+ for (k = 0; k< 100; k++);
}
#pragma acc loop tile(2, 2)
for (i = 1; i < 5; i+=2)
{
- for (j = i + 1; j < 7; j+=i)
+ for (j = i + 1; j < 7; j+=i) // { dg-error "initializer expression" }
{ }
}
#pragma acc loop vector tile(*)
@@ -156,24 +168,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++) // { dg-error "initializer expression" }
+ { }
#pragma acc parallel loop vector tile(*)
for (i = 0; i < 10; i++)
{ }
@@ -227,22 +236,23 @@ kern (void)
#pragma acc loop tile(*, 1)
for (i = 0; i < 10; i++)
{
- for (j = 0; j < 10; i++)
+ for (j = 0; j < 10; i++) /* { dg-error "increment expression" } */
{ }
}
-#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)
for (i = 2; i < 4; i++)
- for (i = 4; i < 6; i++)
+ for (j = 4; j < 6; j++)
+ for (int k = 4; k < 6; k++)
{ }
#pragma acc loop tile(2, 2)
for (i = 1; i < 5; i+=2)
- for (j = i+1; j < 7; i++)
+ for (j = i+1; j < 7; j++) /* { dg-error "initializer expression" } */
{ }
#pragma acc loop vector tile(*)
for (i = 0; i < 10; i++)
@@ -288,22 +298,21 @@ 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++)
{
- for (j = i + 1; j < 7; j += i)
+ for (j = i + 1; j < 7; j += i) /* { dg-error "initializer expression" } */
{ }
}
#pragma acc kernels loop vector tile(*)
Index: testsuite/g++.dg/goacc/template.C
===================================================================
--- testsuite/g++.dg/goacc/template.C (revision 241809)
+++ 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 (revision 0)
+++ testsuite/g++.dg/goacc/tile-1.C (revision 0)
@@ -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
+}
Index: testsuite/gcc.dg/goacc/loop-processing-1.c
===================================================================
--- testsuite/gcc.dg/goacc/loop-processing-1.c (revision 241809)
+++ testsuite/gcc.dg/goacc/loop-processing-1.c (working copy)
@@ -15,4 +15,4 @@ void vector_1 (int *ary, int size)
}
}
-/* { dg-final { scan-tree-dump {OpenACC loops.*Loop 0\(0\).*Loop 14\(1\).*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_HEAD_MARK, 0, 1, 20\);.*Head-0:.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_HEAD_MARK, 0, 1, 20\);.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_FORK, \.data_dep\.[0-9_]+, 0\);.*Tail-0:.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_TAIL_MARK, \.data_dep\.[0-9_]+, 1\);.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_JOIN, \.data_dep\.[0-9_]+, 0\);.*Loop 6\(4\).*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_HEAD_MARK, 0, 1, 6\);.*Head-0:.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_HEAD_MARK, 0, 1, 6\);.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_FORK, \.data_dep\.[0-9_]+, 2\);.*Tail-0:.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_TAIL_MARK, \.data_dep\.[0-9_]+, 1\);.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_JOIN, \.data_dep\.[0-9_]+, 2\);} "oaccdevlow" } } */
+/* { dg-final { scan-tree-dump {OpenACC loops.*Loop 0\(0\).*Loop 24\(1\).*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_HEAD_MARK, 0, 1, 36\);.*Head-0:.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_HEAD_MARK, 0, 1, 36\);.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_FORK, \.data_dep\.[0-9_]+, 0\);.*Tail-0:.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_TAIL_MARK, \.data_dep\.[0-9_]+, 1\);.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_JOIN, \.data_dep\.[0-9_]+, 0\);.*Loop 6\(4\).*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_HEAD_MARK, 0, 2, 4\);.*Head-0:.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_HEAD_MARK, 0, 2, 4\);.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_FORK, \.data_dep\.[0-9_]+, 1\);.*Head-1:.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_HEAD_MARK, \.data_dep\.[0-9_]+, 1\);.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_FORK, \.data_dep\.[0-9_]+, 2\);.*Tail-1:.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_TAIL_MARK, \.data_dep\.[0-9_]+, 2\);.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_JOIN, \.data_dep\.[0-9_]+, 2\);.*Tail-0:.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_TAIL_MARK, \.data_dep\.[0-9_]+, 1\);.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_JOIN, \.data_dep\.[0-9_]+, 1\);} "oaccdevlow" } } */
