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.
@@ -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;
@@ -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:
@@ -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) \
@@ -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 ();
}
@@ -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:
new file mode 100644
@@ -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" } }
new file mode 100644
@@ -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;
+}
@@ -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++)
;
@@ -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)
@@ -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)
@@ -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++)
;