2015-11-05 Cesar Philippidis <cesar@codesourcery.com>
Thomas Schwinge <thomas@codesourcery.com>
James Norris <jnorris@codesourcery.com>
gcc/
* gimplify.c (gimplify_scan_omp_clauses): Add support for
OMP_CLAUSE_TILE. Update handling of OMP_CLAUSE_INDEPENDENT.
(gimplify_adjust_omp_clauses): Likewise.
* omp-low.c (scan_sharing_clauses): Add support for OMP_CLAUSE_TILE.
* tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_TILE.
* tree-pretty-print.c (dump_omp_clause): Handle OMP_CLAUSE_TILE.
* tree.c (omp_clause_num_ops): Add an entry for OMP_CLAUSE_TILE.
(omp_clause_code_name): Likewise.
(walk_tree_1): Handle OMP_CLAUSE_TILE.
* tree.h (OMP_TILE_LIST): New macro.
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.
* c-pragma.h (enum pragma_omp_clause): Add entries for
PRAGMA_OACC_CLAUSE_{INDEPENDENT,TILE,DEFAULT}.
* pt.c (tsubst_omp_clauses): Add support for OMP_CLAUSE_{NUM_GANGS,
NUM_WORKERS,VECTOR_LENGTH,GANG,WORKER,VECTOR,ASYNC,WAIT,TILE,AUTO,
INDEPENDENT,SEQ}.
(tsubst_expr): Add support for OMP_CLAUSE_{KERNELS,PARALLEL,LOOP}.
gcc/c/
* c-parser.c (c_parser_omp_clause_name): Add support for
PRAGMA_OACC_CLAUSE_INDEPENDENT and PRAGMA_OACC_CLAUSE_TILE.
(c_parser_omp_clause_default): Add is_oacc argument. Handle
default(none) in OpenACC.
(c_parser_oacc_shape_clause): Allow pointer variables as gang static
arguments.
(c_parser_oacc_clause_tile): New function.
(c_parser_oacc_all_clauses): Add support for OMP_CLAUSE_DEFAULT,
OMP_CLAUSE_INDEPENDENT and OMP_CLAUSE_TILE.
(OACC_LOOP_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_{PRIVATE,INDEPENDENT,
TILE}.
(OACC_KERNELS_MASK): Add PRAGMA_OACC_CLAUSE_DEFAULT.
(OACC_PARALLEL_MASK): Add PRAGMA_OACC_CLAUSE_{DEFAULT,PRIVATE,
FIRSTPRIVATE}.
(c_parser_omp_all_clauses): Update call to c_parser_omp_clause_default.
(c_parser_oacc_update): Update the error message for missing clauses.
* c-typeck.c (c_finish_omp_clauses): Add support for OMP_CLAUSE_TILE
and OMP_CLAUSE_INDEPENDENT.
gcc/cp/
* parser.c (cp_parser_omp_clause_name): Add support for
PRAGMA_OACC_CLAUSE_INDEPENDENT and PRAGMA_OACC_CLAUSE_TILE.
(cp_parser_oacc_shape_clause): Allow pointer variables as gang static
arguments.
(cp_parser_oacc_clause_tile): New function.
(cp_parser_omp_clause_default): Add is_oacc argument. Handle
default(none) in OpenACC.
(cp_parser_oacc_all_clauses): Add support for
(cp_parser_omp_all_clauses): Update call to
cp_parser_omp_clause_default.
PRAGMA_OACC_CLAUSE_{DEFAULT,INDEPENDENT,TILE,PRIVATE,FIRSTPRIVATE}.
(OACC_LOOP_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_{PRIVATE,INDEPENDENT,
TILE}.
(OACC_KERNELS_MASK): Add PRAGMA_OACC_CLAUSE_DEFAULT.
(OACC_PARALLEL_MASK): Add PRAGMA_OACC_CLAUSE_{DEFAULT,PRIVATE,
FIRSTPRIVATE}.
(cp_parser_oacc_update): Update the error message for missing clauses.
* semantics.c (finish_omp_clauses): Add support for
OMP_CLAUSE_INDEPENDENT and OMP_CLAUSE_TILE.
2015-11-05 Cesar Philippidis <cesar@codesourcery.com>
gcc/fortran/
* openmp.c (gfc_match_omp_clauses): Update support for the tile
and default clauses in OpenACC.
(gfc_match_oacc_update): Error when data clauses are supplied.
(oacc_compatible_clauses): Delete.
(resolve_omp_clauses): Give special care for OpenACC reductions.
Also update error reporting for the tile clause.
(resolve_oacc_loop_blocks): Update error reporting for the tile clause.
* trans-openmp.c (gfc_trans_omp_clauses): Update OMP_CLAUSE_SEQ. Add
OMP_CLAUSE_{AUTO,TILE} and add support the the gang static argument.
(gfc_trans_oacc_combined_directive): Update the list of clauses which
are split to acc loops.
2015-11-05 Cesar Philippidis <cesar@codesourcery.com>
Tom de Vries <tom@codesourcery.com>
Nathan Sidwell <nathan@codesourcery.com>
Thomas Schwinge <thomas@codesourcery.com>
gcc/testsuite/
* c-c++-common/goacc/combined-directives.c: New test.
* c-c++-common/goacc/loop-clauses.c: New test.
* c-c++-common/goacc/tile.c: New test.
* c-c++-common/goacc/loop-shape.c: Add test for pointer variable
as gang static arguments.
* c-c++-common/goacc/update-1.c: Adjust expected error message.
* g++.dg/goacc/template.C: New test.
* gfortran.dg/goacc/combined-directives.f90: New test.
* gfortran.dg/goacc/default.f95: New test.
* gfortran.dg/goacc/default_none.f95: New test.
* gfortran.dg/goacc/firstprivate-1.f95: New test.
* gfortran.dg/goacc/gang-static.f95: New test.
* gfortran.dg/goacc/kernels-loop-inner.f95: New test.
* gfortran.dg/goacc/kernels-loops-adjacent.f95: New test.
* gfortran.dg/goacc/list.f95: Update test.
* gfortran.dg/goacc/loop-2.f95: Likewise.
* gfortran.dg/goacc/loop-4.f95: New test.
* gfortran.dg/goacc/loop-5.f95: New test.
* gfortran.dg/goacc/loop-6.f95: New test.
* gfortran.dg/goacc/loop-tree-1.f90: Update test.
* gfortran.dg/goacc/multi-clause.f90: New test.
* gfortran.dg/goacc/parallel-tree.f95: Update test.
* gfortran.dg/goacc/update.f95: New test.
@@ -867,7 +867,7 @@ c_omp_check_loop_iv_exprs (location_t stmt_loc, tree declv, tree decl,
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)
@@ -876,12 +876,27 @@ 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;
+ /* Reductions belong in both constructs. */
+ case OMP_CLAUSE_REDUCTION:
+ t = copy_node (clauses);
+ OMP_CLAUSE_CHAIN (t) = loop_clauses;
+ loop_clauses = t;
+
+ /* Parallel/kernels clauses. */
default:
OMP_CLAUSE_CHAIN (clauses) = *not_loop_clauses;
*not_loop_clauses = clauses;
@@ -153,6 +153,7 @@ enum pragma_omp_clause {
PRAGMA_OACC_CLAUSE_DEVICEPTR,
PRAGMA_OACC_CLAUSE_GANG,
PRAGMA_OACC_CLAUSE_HOST,
+ PRAGMA_OACC_CLAUSE_INDEPENDENT,
PRAGMA_OACC_CLAUSE_NUM_GANGS,
PRAGMA_OACC_CLAUSE_NUM_WORKERS,
PRAGMA_OACC_CLAUSE_PRESENT,
@@ -162,6 +163,7 @@ enum pragma_omp_clause {
PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE,
PRAGMA_OACC_CLAUSE_SELF,
PRAGMA_OACC_CLAUSE_SEQ,
+ PRAGMA_OACC_CLAUSE_TILE,
PRAGMA_OACC_CLAUSE_VECTOR,
PRAGMA_OACC_CLAUSE_VECTOR_LENGTH,
PRAGMA_OACC_CLAUSE_WAIT,
@@ -169,6 +171,7 @@ enum pragma_omp_clause {
PRAGMA_OACC_CLAUSE_COLLAPSE = PRAGMA_OMP_CLAUSE_COLLAPSE,
PRAGMA_OACC_CLAUSE_COPYIN = PRAGMA_OMP_CLAUSE_COPYIN,
PRAGMA_OACC_CLAUSE_DEVICE = PRAGMA_OMP_CLAUSE_DEVICE,
+ PRAGMA_OACC_CLAUSE_DEFAULT = PRAGMA_OMP_CLAUSE_DEFAULT,
PRAGMA_OACC_CLAUSE_FIRSTPRIVATE = PRAGMA_OMP_CLAUSE_FIRSTPRIVATE,
PRAGMA_OACC_CLAUSE_IF = PRAGMA_OMP_CLAUSE_IF,
PRAGMA_OACC_CLAUSE_PRIVATE = PRAGMA_OMP_CLAUSE_PRIVATE,
@@ -10008,6 +10008,8 @@ c_parser_omp_clause_name (c_parser *parser)
case 'i':
if (!strcmp ("inbranch", p))
result = PRAGMA_OMP_CLAUSE_INBRANCH;
+ else if (!strcmp ("independent", p))
+ result = PRAGMA_OACC_CLAUSE_INDEPENDENT;
else if (!strcmp ("is_device_ptr", p))
result = PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR;
break;
@@ -10104,6 +10106,8 @@ c_parser_omp_clause_name (c_parser *parser)
result = PRAGMA_OMP_CLAUSE_THREAD_LIMIT;
else if (!strcmp ("threads", p))
result = PRAGMA_OMP_CLAUSE_THREADS;
+ else if (!strcmp ("tile", p))
+ result = PRAGMA_OACC_CLAUSE_TILE;
else if (!strcmp ("to", p))
result = PRAGMA_OMP_CLAUSE_TO;
break;
@@ -10541,10 +10545,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)
+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;
@@ -10565,7 +10572,7 @@ c_parser_omp_clause_default (c_parser *parser, tree list)
break;
case 's':
- if (strcmp ("shared", p) != 0)
+ if (strcmp ("shared", p) != 0 || is_oacc)
goto invalid_kind;
kind = OMP_CLAUSE_DEFAULT_SHARED;
break;
@@ -10579,7 +10586,10 @@ c_parser_omp_clause_default (c_parser *parser, tree list)
else
{
invalid_kind:
- c_parser_error (parser, "expected %<none%> or %<shared%>");
+ if (is_oacc)
+ c_parser_error (parser, "expected %<none%>");
+ else
+ c_parser_error (parser, "expected %<none%> or %<shared%>");
}
c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, "expected %<)%>");
@@ -11236,7 +11246,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;
@@ -11378,6 +11391,79 @@ c_parser_oacc_clause_async (c_parser *parser, tree list)
return list;
}
+/* OpenACC 2.0:
+ tile ( size-expr-list ) */
+
+static tree
+c_parser_oacc_clause_tile (c_parser *parser, tree list)
+{
+ tree c, expr = error_mark_node;
+ location_t loc, expr_loc;
+ tree tile = NULL_TREE;
+
+ 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 %<(%>"))
+ return list;
+
+ do
+ {
+ 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;
+ }
+ else
+ {
+ expr_loc = c_parser_peek_token (parser)->location;
+ expr = c_parser_expr_no_commas (parser, NULL).value;
+
+ if (expr == error_mark_node)
+ {
+ c_parser_skip_until_found (parser, CPP_CLOSE_PAREN,
+ "expected %<)%>");
+ 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)
+ {
+ warning_at (expr_loc, 0,"%<tile%> value must be positive");
+ expr = integer_one_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));
+
+ /* Consume the trailing ')'. */
+ c_parser_consume_token (parser);
+
+ c = build_omp_clause (loc, OMP_CLAUSE_TILE);
+ tile = nreverse (tile);
+ OMP_CLAUSE_TILE_LIST (c) = tile;
+ OMP_CLAUSE_CHAIN (c) = list;
+ return c;
+}
+
/* OpenACC:
wait ( int-expr-list ) */
@@ -12605,6 +12691,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "delete";
break;
+ case PRAGMA_OMP_CLAUSE_DEFAULT:
+ clauses = c_parser_omp_clause_default (parser, clauses, true);
+ c_name = "default";
+ break;
case PRAGMA_OACC_CLAUSE_DEVICE:
clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "device";
@@ -12630,6 +12720,11 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
clauses = c_parser_omp_clause_if (parser, clauses, false);
c_name = "if";
break;
+ case PRAGMA_OACC_CLAUSE_INDEPENDENT:
+ clauses = c_parser_oacc_simple_clause (parser, OMP_CLAUSE_INDEPENDENT,
+ clauses);
+ c_name = "independent";
+ break;
case PRAGMA_OACC_CLAUSE_NUM_GANGS:
clauses = c_parser_omp_clause_num_gangs (parser, clauses);
c_name = "num_gangs";
@@ -12675,6 +12770,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
clauses);
c_name = "seq";
break;
+ case PRAGMA_OACC_CLAUSE_TILE:
+ clauses = c_parser_oacc_clause_tile (parser, clauses);
+ c_name = "tile";
+ break;
case PRAGMA_OACC_CLAUSE_VECTOR:
c_name = "vector";
clauses = c_parser_oacc_shape_clause (parser, OMP_CLAUSE_VECTOR,
@@ -12756,7 +12855,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:
@@ -13169,13 +13268,15 @@ c_parser_oacc_enter_exit_data (c_parser *parser, bool enter)
#define OACC_LOOP_CLAUSE_MASK \
( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COLLAPSE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRIVATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_GANG) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WORKER) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_AUTO) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_INDEPENDENT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) )
-
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_TILE) )
static tree
c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
omp_clause_mask mask, tree *cclauses)
@@ -13220,6 +13321,7 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
@@ -13235,8 +13337,11 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRIVATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FIRSTPRIVATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
@@ -13318,7 +13423,7 @@ c_parser_oacc_update (c_parser *parser)
{
error_at (loc,
"%<#pragma acc update%> must contain at least one "
- "%<device%> or %<host/self%> clause");
+ "%<device%> or %<host%> or %<self%> clause");
return;
}
@@ -13011,10 +13011,12 @@ c_finish_omp_clauses (tree clauses, bool is_omp, bool declare_simd)
case OMP_CLAUSE_ASYNC:
case OMP_CLAUSE_WAIT:
case OMP_CLAUSE_AUTO:
+ case OMP_CLAUSE_INDEPENDENT:
case OMP_CLAUSE_SEQ:
case OMP_CLAUSE_GANG:
case OMP_CLAUSE_WORKER:
case OMP_CLAUSE_VECTOR:
+ case OMP_CLAUSE_TILE:
pc = &OMP_CLAUSE_CHAIN (c);
continue;
@@ -29125,6 +29125,8 @@ cp_parser_omp_clause_name (cp_parser *parser)
case 'i':
if (!strcmp ("inbranch", p))
result = PRAGMA_OMP_CLAUSE_INBRANCH;
+ else if (!strcmp ("independent", p))
+ result = PRAGMA_OACC_CLAUSE_INDEPENDENT;
else if (!strcmp ("is_device_ptr", p))
result = PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR;
break;
@@ -29219,6 +29221,8 @@ cp_parser_omp_clause_name (cp_parser *parser)
result = PRAGMA_OMP_CLAUSE_THREAD_LIMIT;
else if (!strcmp ("threads", p))
result = PRAGMA_OMP_CLAUSE_THREADS;
+ else if (!strcmp ("tile", p))
+ result = PRAGMA_OACC_CLAUSE_TILE;
else if (!strcmp ("to", p))
result = PRAGMA_OMP_CLAUSE_TO;
break;
@@ -29684,7 +29688,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;
@@ -29752,6 +29759,52 @@ cp_parser_oacc_shape_clause (cp_parser *parser, omp_clause_code kind,
return list;
}
+/* OpenACC 2.0:
+ tile ( size-expr-list ) */
+
+static tree
+cp_parser_oacc_clause_tile (cp_parser *parser, location_t clause_loc, tree list)
+{
+ tree c, expr = error_mark_node;
+ tree tile = NULL_TREE;
+
+ check_no_duplicate_clause (list, OMP_CLAUSE_TILE, "tile", clause_loc);
+
+ if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
+ return list;
+
+ do
+ {
+ 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;
+ }
+ else
+ expr = cp_parser_assignment_expression (parser, NULL, false, false);
+
+ 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));
+
+ /* Consume the trailing ')'. */
+ cp_lexer_consume_token (parser->lexer);
+
+ c = build_omp_clause (clause_loc, OMP_CLAUSE_TILE);
+ tile = nreverse (tile);
+ OMP_CLAUSE_TILE_LIST (c) = tile;
+ OMP_CLAUSE_CHAIN (c) = list;
+ return c;
+}
+
/* OpenACC 2.0
Parse wait clause or directive parameters. */
@@ -29859,10 +29912,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)
+cp_parser_omp_clause_default (cp_parser *parser, tree list,
+ location_t location, bool is_oacc)
{
enum omp_clause_default_kind kind = OMP_CLAUSE_DEFAULT_UNSPECIFIED;
tree c;
@@ -29883,7 +29940,7 @@ cp_parser_omp_clause_default (cp_parser *parser, tree list, location_t location)
break;
case 's':
- if (strcmp ("shared", p) != 0)
+ if (strcmp ("shared", p) != 0 || is_oacc)
goto invalid_kind;
kind = OMP_CLAUSE_DEFAULT_SHARED;
break;
@@ -29897,7 +29954,10 @@ cp_parser_omp_clause_default (cp_parser *parser, tree list, location_t location)
else
{
invalid_kind:
- cp_parser_error (parser, "expected %<none%> or %<shared%>");
+ 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))
@@ -31444,6 +31504,10 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "delete";
break;
+ case PRAGMA_OMP_CLAUSE_DEFAULT:
+ clauses = cp_parser_omp_clause_default (parser, clauses, here, true);
+ c_name = "default";
+ break;
case PRAGMA_OACC_CLAUSE_DEVICE:
clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "device";
@@ -31452,6 +31516,11 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
clauses = cp_parser_oacc_data_clause_deviceptr (parser, clauses);
c_name = "deviceptr";
break;
+ case PRAGMA_OACC_CLAUSE_FIRSTPRIVATE:
+ clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_FIRSTPRIVATE,
+ clauses);
+ c_name = "firstprivate";
+ break;
case PRAGMA_OACC_CLAUSE_GANG:
c_name = "gang";
clauses = cp_parser_oacc_shape_clause (parser, OMP_CLAUSE_GANG,
@@ -31465,6 +31534,12 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
clauses = cp_parser_omp_clause_if (parser, clauses, here, false);
c_name = "if";
break;
+ case PRAGMA_OACC_CLAUSE_INDEPENDENT:
+ clauses = cp_parser_oacc_simple_clause (parser,
+ OMP_CLAUSE_INDEPENDENT,
+ clauses, here);
+ c_name = "independent";
+ break;
case PRAGMA_OACC_CLAUSE_NUM_GANGS:
code = OMP_CLAUSE_NUM_GANGS;
c_name = "num_gangs";
@@ -31497,6 +31572,11 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "present_or_create";
break;
+ case PRAGMA_OACC_CLAUSE_PRIVATE:
+ clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_PRIVATE,
+ clauses);
+ c_name = "private";
+ break;
case PRAGMA_OACC_CLAUSE_REDUCTION:
clauses = cp_parser_omp_clause_reduction (parser, clauses);
c_name = "reduction";
@@ -31510,6 +31590,10 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
clauses, here);
c_name = "seq";
break;
+ case PRAGMA_OACC_CLAUSE_TILE:
+ clauses = cp_parser_oacc_clause_tile (parser, here, clauses);
+ c_name = "tile";
+ break;
case PRAGMA_OACC_CLAUSE_VECTOR:
c_name = "vector";
clauses = cp_parser_oacc_shape_clause (parser, OMP_CLAUSE_VECTOR,
@@ -31598,7 +31682,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);
+ token->location, false);
c_name = "default";
break;
case PRAGMA_OMP_CLAUSE_FINAL:
@@ -34493,12 +34577,15 @@ cp_parser_oacc_enter_exit_data (cp_parser *parser, cp_token *pragma_tok,
#define OACC_LOOP_CLAUSE_MASK \
( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COLLAPSE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRIVATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_GANG) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WORKER) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_AUTO) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_INDEPENDENT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) )
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_TILE))
static tree
cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
@@ -34543,6 +34630,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
@@ -34558,7 +34646,9 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FIRSTPRIVATE) \
| (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) \
@@ -34567,6 +34657,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRIVATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
@@ -34642,7 +34733,7 @@ cp_parser_oacc_update (cp_parser *parser, cp_token *pragma_tok)
{
error_at (pragma_tok->location,
"%<#pragma acc update%> must contain at least one "
- "%<device%> or %<host/self%> clause");
+ "%<device%> or %<host%> or %<self%> clause");
return NULL_TREE;
}
@@ -14395,6 +14395,13 @@ tsubst_omp_clauses (tree clauses, bool declare_simd, bool allow_fields,
case OMP_CLAUSE_PRIORITY:
case OMP_CLAUSE_ORDERED:
case OMP_CLAUSE_HINT:
+ case OMP_CLAUSE_NUM_GANGS:
+ case OMP_CLAUSE_NUM_WORKERS:
+ case OMP_CLAUSE_VECTOR_LENGTH:
+ case OMP_CLAUSE_WORKER:
+ case OMP_CLAUSE_VECTOR:
+ 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);
@@ -14419,6 +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_GANG:
case OMP_CLAUSE_ALIGNED:
OMP_CLAUSE_DECL (nc)
= tsubst_omp_clause_decl (OMP_CLAUSE_DECL (oc), args, complain,
@@ -14461,6 +14469,22 @@ tsubst_omp_clauses (tree clauses, bool declare_simd, bool allow_fields,
case OMP_CLAUSE_THREADS:
case OMP_CLAUSE_SIMD:
case OMP_CLAUSE_DEFAULTMAP:
+ case OMP_CLAUSE_INDEPENDENT:
+ 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 ();
@@ -15231,6 +15255,15 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
}
break;
+ case OACC_KERNELS:
+ case OACC_PARALLEL:
+ tmp = tsubst_omp_clauses (OMP_CLAUSES (t), false, false, args, complain,
+ in_decl);
+ stmt = begin_omp_parallel ();
+ RECUR (OMP_BODY (t));
+ finish_omp_construct (TREE_CODE (t), stmt, tmp);
+ break;
+
case OMP_PARALLEL:
r = push_omp_privatization_clauses (OMP_PARALLEL_COMBINED (t));
tmp = tsubst_omp_clauses (OMP_PARALLEL_CLAUSES (t), false, true,
@@ -15261,6 +15294,7 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
case CILK_FOR:
case OMP_DISTRIBUTE:
case OMP_TASKLOOP:
+ case OACC_LOOP:
{
tree clauses, body, pre_body;
tree declv = NULL_TREE, initv = NULL_TREE, condv = NULL_TREE;
@@ -15269,7 +15303,8 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
int i;
r = push_omp_privatization_clauses (OMP_FOR_INIT (t) == NULL_TREE);
- clauses = tsubst_omp_clauses (OMP_FOR_CLAUSES (t), false, true,
+ clauses = tsubst_omp_clauses (OMP_FOR_CLAUSES (t), false,
+ TREE_CODE (t) != OACC_LOOP,
args, complain, in_decl);
if (OMP_FOR_INIT (t) != NULL_TREE)
{
@@ -15339,9 +15374,11 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
pop_omp_privatization_clauses (r);
break;
+ case OACC_DATA:
case OMP_TARGET_DATA:
case OMP_TARGET:
- tmp = tsubst_omp_clauses (OMP_CLAUSES (t), false, true,
+ tmp = tsubst_omp_clauses (OMP_CLAUSES (t), false,
+ TREE_CODE (t) != OACC_DATA,
args, complain, in_decl);
keep_next_level (true);
stmt = begin_omp_structured_block ();
@@ -15395,6 +15432,16 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
add_stmt (t);
break;
+ case OACC_ENTER_DATA:
+ case OACC_EXIT_DATA:
+ case OACC_UPDATE:
+ tmp = tsubst_omp_clauses (OMP_STANDALONE_CLAUSES (t), false, false,
+ args, complain, in_decl);
+ t = copy_node (t);
+ OMP_STANDALONE_CLAUSES (t) = tmp;
+ add_stmt (t);
+ break;
+
case OMP_ORDERED:
tmp = tsubst_omp_clauses (OMP_ORDERED_CLAUSES (t), false, true,
args, complain, in_decl);
@@ -6855,9 +6855,47 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd)
case OMP_CLAUSE_DEFAULTMAP:
case OMP_CLAUSE__CILK_FOR_COUNT_:
case OMP_CLAUSE_AUTO:
+ case OMP_CLAUSE_INDEPENDENT:
case OMP_CLAUSE_SEQ:
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_ORDERED:
ordered_seen = true;
break;
@@ -703,6 +703,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
OMP_MAP_FORCE_FROM))
continue;
if ((mask & OMP_CLAUSE_TILE)
+ && !c->tile_list
&& match_oacc_expr_list ("tile (", &c->tile_list, true) == MATCH_YES)
continue;
if ((mask & OMP_CLAUSE_SEQ) && !c->seq
@@ -856,12 +857,14 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
if ((mask & OMP_CLAUSE_DEFAULT)
&& c->default_sharing == OMP_DEFAULT_UNKNOWN)
{
- if (gfc_match ("default ( shared )") == MATCH_YES)
+ if (gfc_match ("default ( none )") == MATCH_YES)
+ c->default_sharing = OMP_DEFAULT_NONE;
+ else if (openacc)
+ /* c->default_sharing = OMP_DEFAULT_UNKNOWN */;
+ else if (gfc_match ("default ( shared )") == MATCH_YES)
c->default_sharing = OMP_DEFAULT_SHARED;
else if (gfc_match ("default ( private )") == MATCH_YES)
c->default_sharing = OMP_DEFAULT_PRIVATE;
- else if (gfc_match ("default ( none )") == MATCH_YES)
- c->default_sharing = OMP_DEFAULT_NONE;
else if (gfc_match ("default ( firstprivate )") == MATCH_YES)
c->default_sharing = OMP_DEFAULT_FIRSTPRIVATE;
if (c->default_sharing != OMP_DEFAULT_UNKNOWN)
@@ -1304,10 +1307,19 @@ match
gfc_match_oacc_update (void)
{
gfc_omp_clauses *c;
+ locus here = gfc_current_locus;
+
if (gfc_match_omp_clauses (&c, OACC_UPDATE_CLAUSES, false, false, true)
!= MATCH_YES)
return MATCH_ERROR;
+ if (!c->lists[OMP_LIST_MAP])
+ {
+ gfc_error ("%<acc update%> must contain at least one "
+ "%<device%> or %<host%> or %<self%> clause at %L", &here);
+ return MATCH_ERROR;
+ }
+
new_st.op = EXEC_OACC_UPDATE;
new_st.ext.omp_clauses = c;
return MATCH_YES;
@@ -2846,30 +2858,6 @@ resolve_omp_udr_clause (gfc_omp_namelist *n, gfc_namespace *ns,
return copy;
}
-/* Returns true if clause in list 'list' is compatible with any of
- of the clauses in lists [0..list-1]. E.g., a reduction variable may
- appear in both reduction and private clauses, so this function
- will return true in this case. */
-
-static bool
-oacc_compatible_clauses (gfc_omp_clauses *clauses, int list,
- gfc_symbol *sym, bool openacc)
-{
- gfc_omp_namelist *n;
-
- if (!openacc)
- return false;
-
- if (list != OMP_LIST_REDUCTION)
- return false;
-
- for (n = clauses->lists[OMP_LIST_FIRST]; n; n = n->next)
- if (n->sym == sym)
- return true;
-
- return false;
-}
-
/* OpenMP directive resolving routines. */
static void
@@ -2975,11 +2963,11 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
&& list != OMP_LIST_DEPEND
&& (list != OMP_LIST_MAP || openacc)
&& list != OMP_LIST_FROM
- && list != OMP_LIST_TO)
+ && list != OMP_LIST_TO
+ && (list != OMP_LIST_REDUCTION || !openacc))
for (n = omp_clauses->lists[list]; n; n = n->next)
{
- if (n->sym->mark && !oacc_compatible_clauses (omp_clauses, list,
- n->sym, openacc))
+ if (n->sym->mark)
gfc_error ("Symbol %qs present on multiple clauses at %L",
n->sym->name, &n->where);
else
@@ -3028,6 +3016,22 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
n->sym->mark = 1;
}
+ /* OpenACC reductions. */
+ if (openacc)
+ {
+ for (n = omp_clauses->lists[OMP_LIST_REDUCTION]; n; n = n->next)
+ n->sym->mark = 0;
+
+ for (n = omp_clauses->lists[OMP_LIST_REDUCTION]; n; n = n->next)
+ {
+ if (n->sym->mark)
+ gfc_error ("Symbol %qs present on multiple clauses at %L",
+ n->sym->name, &n->where);
+ else
+ n->sym->mark = 1;
+ }
+ }
+
for (n = omp_clauses->lists[OMP_LIST_TO]; n; n = n->next)
n->sym->mark = 0;
for (n = omp_clauses->lists[OMP_LIST_FROM]; n; n = n->next)
@@ -4528,22 +4532,8 @@ resolve_oacc_loop_blocks (gfc_code *code)
if (code->ext.omp_clauses->vector)
gfc_error ("Clause AUTO conflicts with VECTOR at %L", &code->loc);
}
- if (!code->ext.omp_clauses->tile_list)
- {
- if (code->ext.omp_clauses->gang)
- {
- if (code->ext.omp_clauses->worker)
- gfc_error ("Clause GANG conflicts with WORKER at %L", &code->loc);
- if (code->ext.omp_clauses->vector)
- gfc_error ("Clause GANG conflicts with VECTOR at %L", &code->loc);
- }
- if (code->ext.omp_clauses->worker)
- if (code->ext.omp_clauses->vector)
- gfc_error ("Clause WORKER conflicts with VECTOR at %L", &code->loc);
- }
- else if (code->ext.omp_clauses->gang
- && code->ext.omp_clauses->worker
- && code->ext.omp_clauses->vector)
+ if (code->ext.omp_clauses->tile_list && code->ext.omp_clauses->gang
+ && code->ext.omp_clauses->worker && code->ext.omp_clauses->vector)
gfc_error ("Tiled loop cannot be parallelized across gangs, workers and "
"vectors at the same time at %L", &code->loc);
@@ -4564,10 +4554,21 @@ resolve_oacc_loop_blocks (gfc_code *code)
{
num++;
if (el->expr == NULL)
- continue;
- resolve_oacc_positive_int_expr (el->expr, "TILE");
- if (el->expr->expr_type != EXPR_CONSTANT)
- gfc_error ("TILE requires constant expression at %L", &code->loc);
+ {
+ /* NULL expressions are used to represent '*' arguments.
+ Convert those to a -1 expressions. */
+ el->expr = gfc_get_constant_expr (BT_INTEGER,
+ gfc_default_integer_kind,
+ &code->loc);
+ mpz_set_si (el->expr->value.integer, -1);
+ }
+ else
+ {
+ resolve_oacc_positive_int_expr (el->expr, "TILE");
+ if (el->expr->expr_type != EXPR_CONSTANT)
+ gfc_error ("TILE requires constant expression at %L",
+ &code->loc);
+ }
}
resolve_oacc_nested_loops (code, code->block->next, num, "tiled");
}
@@ -2534,8 +2534,12 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
}
if (clauses->seq)
{
- c = build_omp_clause (where.lb->location, OMP_CLAUSE_ORDERED);
- OMP_CLAUSE_ORDERED_EXPR (c) = NULL_TREE;
+ c = build_omp_clause (where.lb->location, OMP_CLAUSE_SEQ);
+ omp_clauses = gfc_trans_add_clause (c, omp_clauses);
+ }
+ if (clauses->par_auto)
+ {
+ c = build_omp_clause (where.lb->location, OMP_CLAUSE_AUTO);
omp_clauses = gfc_trans_add_clause (c, omp_clauses);
}
if (clauses->independent)
@@ -2579,6 +2583,21 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
OMP_CLAUSE_VECTOR_LENGTH_EXPR (c) = vector_length_var;
omp_clauses = gfc_trans_add_clause (c, omp_clauses);
}
+ if (clauses->tile_list)
+ {
+ vec<tree, va_gc> *tvec;
+ gfc_expr_list *el;
+
+ vec_alloc (tvec, 4);
+
+ for (el = clauses->tile_list; el; el = el->next)
+ vec_safe_push (tvec, gfc_convert_expr_to_tree (block, el->expr));
+
+ c = build_omp_clause (where.lb->location, OMP_CLAUSE_TILE);
+ OMP_CLAUSE_TILE_LIST (c) = build_tree_list_vec (tvec);
+ omp_clauses = gfc_trans_add_clause (c, omp_clauses);
+ tvec->truncate (0);
+ }
if (clauses->vector)
{
if (clauses->vector_expr)
@@ -2618,7 +2637,17 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
tree gang_var
= gfc_convert_expr_to_tree (block, clauses->gang_expr);
c = build_omp_clause (where.lb->location, OMP_CLAUSE_GANG);
- OMP_CLAUSE_GANG_EXPR (c) = gang_var;
+ if (clauses->gang_static)
+ OMP_CLAUSE_GANG_STATIC_EXPR (c) = gang_var;
+ else
+ OMP_CLAUSE_GANG_EXPR (c) = gang_var;
+ omp_clauses = gfc_trans_add_clause (c, omp_clauses);
+ }
+ else if (clauses->gang_static)
+ {
+ /* This corresponds to gang (static: *). */
+ c = build_omp_clause (where.lb->location, OMP_CLAUSE_GANG);
+ OMP_CLAUSE_GANG_STATIC_EXPR (c) = integer_minus_one_node;
omp_clauses = gfc_trans_add_clause (c, omp_clauses);
}
else
@@ -3449,16 +3478,33 @@ gfc_trans_oacc_combined_directive (gfc_code *code)
sizeof (construct_clauses));
loop_clauses.collapse = construct_clauses.collapse;
loop_clauses.gang = construct_clauses.gang;
+ loop_clauses.gang_expr = construct_clauses.gang_expr;
+ loop_clauses.gang_static = construct_clauses.gang_static;
loop_clauses.vector = construct_clauses.vector;
+ loop_clauses.vector_expr = construct_clauses.vector_expr;
loop_clauses.worker = construct_clauses.worker;
+ loop_clauses.worker_expr = construct_clauses.worker_expr;
loop_clauses.seq = construct_clauses.seq;
+ loop_clauses.par_auto = construct_clauses.par_auto;
loop_clauses.independent = construct_clauses.independent;
- construct_clauses.collapse = 0;
+ loop_clauses.tile_list = construct_clauses.tile_list;
+ loop_clauses.lists[OMP_LIST_PRIVATE]
+ = construct_clauses.lists[OMP_LIST_PRIVATE];
+ loop_clauses.lists[OMP_LIST_REDUCTION]
+ = construct_clauses.lists[OMP_LIST_REDUCTION];
construct_clauses.gang = false;
+ construct_clauses.gang_expr = NULL;
+ construct_clauses.gang_static = false;
construct_clauses.vector = false;
+ construct_clauses.vector_expr = NULL;
construct_clauses.worker = false;
+ construct_clauses.worker_expr = NULL;
construct_clauses.seq = false;
+ construct_clauses.par_auto = false;
+ construct_clauses.independent = false;
construct_clauses.independent = false;
+ construct_clauses.tile_list = NULL;
+ construct_clauses.lists[OMP_LIST_PRIVATE] = NULL;
oacc_clauses = gfc_trans_omp_clauses (&block, &construct_clauses,
code->loc);
}
@@ -7141,9 +7141,18 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
remove = true;
break;
+ case OMP_CLAUSE_TILE:
+ for (tree list = OMP_CLAUSE_TILE_LIST (c); !remove && list;
+ list = TREE_CHAIN (list))
+ {
+ if (gimplify_expr (&TREE_VALUE (list), pre_p, NULL,
+ is_gimple_val, fb_rvalue) == GS_ERROR)
+ remove = true;
+ }
+ break;
+
case OMP_CLAUSE_DEVICE_RESIDENT:
case OMP_CLAUSE_USE_DEVICE:
- case OMP_CLAUSE_INDEPENDENT:
remove = true;
break;
@@ -7153,6 +7162,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
case OMP_CLAUSE_COLLAPSE:
case OMP_CLAUSE_AUTO:
case OMP_CLAUSE_SEQ:
+ case OMP_CLAUSE_INDEPENDENT:
case OMP_CLAUSE_MERGEABLE:
case OMP_CLAUSE_PROC_BIND:
case OMP_CLAUSE_SAFELEN:
@@ -7663,6 +7673,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, tree *list_p,
case OMP_CLAUSE_VECTOR:
case OMP_CLAUSE_AUTO:
case OMP_CLAUSE_SEQ:
+ case OMP_CLAUSE_TILE:
break;
default:
@@ -2129,6 +2129,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_GANG:
case OMP_CLAUSE_WORKER:
case OMP_CLAUSE_VECTOR:
+ case OMP_CLAUSE_TILE:
break;
case OMP_CLAUSE_ALIGNED:
@@ -2309,6 +2310,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_GANG:
case OMP_CLAUSE_WORKER:
case OMP_CLAUSE_VECTOR:
+ case OMP_CLAUSE_TILE:
break;
case OMP_CLAUSE_DEVICE_RESIDENT:
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++)
;
new file mode 100644
@@ -0,0 +1,73 @@
+/* { dg-do compile } */
+
+int
+main ()
+{
+ int i, *a, b;
+
+#pragma acc parallel loop tile (10)
+ for (i = 0; i < 100; i++)
+ ;
+
+#pragma acc parallel loop tile (*)
+ for (i = 0; i < 100; i++)
+ ;
+
+#pragma acc parallel loop tile (10, *)
+ for (i = 0; i < 100; i++)
+ ;
+
+#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++)
+ ;
+
+ return 0;
+}
@@ -1,7 +1,7 @@
void
f (void)
{
-#pragma acc update /* { dg-error "'#pragma acc update' must contain at least one 'device' or 'host/self' clause" } */
+#pragma acc update /* { dg-error "'#pragma acc update' must contain at least one 'device' or 'host' or 'self' clause" } */
int i = 0;
int a[10];
new file mode 100644
@@ -0,0 +1,141 @@
+// 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" }
+
+#pragma acc routine
+template <typename T> T
+accDouble(int val)
+{
+ return val * 2;
+}
+
+template<typename T> T
+oacc_parallel_copy (T a)
+{
+ T b = 0;
+ char w = 1;
+ int x = 2;
+ float y = 3;
+ double z = 4;
+
+#pragma acc parallel num_gangs (a) num_workers (a) vector_length (a) default (none) copyout (b) copyin (a)
+ {
+ b = a;
+ }
+
+#pragma acc parallel num_gangs (a) copy (w, x, y, z)
+ {
+ w = accDouble<char>(w);
+ x = accDouble<int>(x);
+ y = accDouble<float>(y);
+ z = accDouble<double>(z);
+ }
+
+#pragma acc parallel num_gangs (a) if (1)
+ {
+#pragma acc loop auto tile (a, 3)
+ for (int i = 0; i < a; i++)
+ for (int j = 0; j < 5; j++)
+ b = a;
+
+#pragma acc loop seq
+ for (int i = 0; i < a; i++)
+ b = a;
+ }
+
+ T c;
+
+#pragma acc parallel num_workers (10)
+ {
+#pragma acc atomic capture
+ c = b++;
+
+#pragma atomic update
+ c++;
+
+#pragma acc atomic read
+ b = a;
+
+#pragma acc atomic write
+ b = a;
+ }
+
+//#pragma acc parallel reduction (+:c)
+// {
+// c = 1;
+// }
+
+#pragma acc data if (1) copy (b)
+ {
+ #pragma acc parallel
+ {
+ b = a;
+ }
+ }
+
+#pragma acc enter data copyin (b)
+#pragma acc parallel present (b)
+ {
+ b = a;
+ }
+
+#pragma acc update host (b)
+#pragma acc update self (b)
+#pragma acc update device (b)
+#pragma acc exit data delete (b)
+
+ return b;
+}
+
+template<typename T> T
+oacc_kernels_copy (T a)
+{
+ T b = 0;
+ T c = 0;
+ char w = 1;
+ int x = 2;
+ float y = 3;
+ double z = 4;
+
+#pragma acc kernels copy (w, x, y, z)
+ {
+ w = accDouble<char>(w);
+ x = accDouble<int>(x);
+ y = accDouble<float>(y);
+ z = accDouble<double>(z);
+ }
+
+#pragma acc kernels copyout (b) copyin (a)
+ b = a;
+
+//#pragma acc kernels loop reduction (+:c)
+// for (int i = 0; i < 10; i++)
+// {
+// c = 1;
+// }
+
+#pragma acc data if (1) copy (b)
+ {
+ #pragma acc kernels
+ {
+ b = a;
+ }
+ }
+
+#pragma acc enter data copyin (b)
+#pragma acc kernels present (b)
+ {
+ b = a;
+ }
+ return b;
+}
+
+int
+main ()
+{
+ int b = oacc_parallel_copy<int> (5);
+ int c = oacc_kernels_copy<int> (5);
+
+ return b + c;
+}
new file mode 100644
@@ -0,0 +1,155 @@
+! Exercise combined OpenACC directives.
+
+! { 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.
+! { dg-prune-output "sorry, unimplemented" }
+
+! Update the reduction tests.
+
+subroutine test
+ implicit none
+ integer a(100), i, j, z
+
+ ! PARALLEL
+
+ !$acc parallel loop collapse (2)
+ do i = 1, 100
+ do j = 1, 10
+ end do
+ end do
+ !$acc end parallel loop
+
+ !$acc parallel loop gang
+ do i = 1, 100
+ end do
+ !$acc end parallel loop
+
+ !$acc parallel loop worker
+ do i = 1, 100
+ do j = 1, 10
+ end do
+ end do
+ !$acc end parallel loop
+
+ !$acc parallel loop vector
+ do i = 1, 100
+ do j = 1, 10
+ end do
+ end do
+ !$acc end parallel loop
+
+ !$acc parallel loop seq
+ do i = 1, 100
+ do j = 1, 10
+ end do
+ end do
+ !$acc end parallel loop
+
+ !$acc parallel loop auto
+ do i = 1, 100
+ do j = 1, 10
+ end do
+ end do
+ !$acc end parallel loop
+
+ !$acc parallel loop tile (2, 3)
+ do i = 1, 100
+ do j = 1, 10
+ end do
+ end do
+ !$acc end parallel loop
+
+ !$acc parallel loop independent
+ do i = 1, 100
+ end do
+ !$acc end parallel loop
+
+ !$acc parallel loop private (z)
+ do i = 1, 100
+ z = 0
+ end do
+ !$acc end parallel loop
+
+! !$acc parallel loop reduction (+:z) copy (z)
+! do i = 1, 100
+! end do
+! !$acc end parallel loop
+
+ ! KERNELS
+
+ !$acc kernels loop collapse (2)
+ do i = 1, 100
+ do j = 1, 10
+ end do
+ end do
+ !$acc end kernels loop
+
+ !$acc kernels loop gang
+ do i = 1, 100
+ end do
+ !$acc end kernels loop
+
+ !$acc kernels loop worker
+ do i = 1, 100
+ do j = 1, 10
+ end do
+ end do
+ !$acc end kernels loop
+
+ !$acc kernels loop vector
+ do i = 1, 100
+ do j = 1, 10
+ end do
+ end do
+ !$acc end kernels loop
+
+ !$acc kernels loop seq
+ do i = 1, 100
+ do j = 1, 10
+ end do
+ end do
+ !$acc end kernels loop
+
+ !$acc kernels loop auto
+ do i = 1, 100
+ do j = 1, 10
+ end do
+ end do
+ !$acc end kernels loop
+
+ !$acc kernels loop tile (2, 3)
+ do i = 1, 100
+ do j = 1, 10
+ end do
+ end do
+ !$acc end kernels loop
+
+ !$acc kernels loop independent
+ do i = 1, 100
+ end do
+ !$acc end kernels loop
+
+ !$acc kernels loop private (z)
+ do i = 1, 100
+ z = 0
+ end do
+ !$acc end kernels loop
+
+! !$acc kernels loop reduction (+:z) copy (z)
+! do i = 1, 100
+! end do
+! !$acc end kernels loop
+end subroutine test
+
+! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. collapse.2." 2 "gimple" } }
+! { dg-final { scan-tree-dump-times "acc loop private.i. gang" 2 "gimple" } }
+! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. worker" 2 "gimple" } }
+! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. vector" 2 "gimple" } }
+! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. seq" 2 "gimple" } }
+! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. auto" 2 "gimple" } }
+! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. tile.2, 3" 2 "gimple" } }
+! { dg-final { scan-tree-dump-times "acc loop private.i. independent" 2 "gimple" } }
+! { dg-final { scan-tree-dump-times "private.z" 2 "gimple" } }
new file mode 100644
@@ -0,0 +1,17 @@
+! { dg-do compile }
+
+program tile
+ integer i, j, a
+
+ !$acc parallel default (shared) ! { dg-error "Unclassifiable OpenACC directive" }
+ !$acc end parallel ! { dg-error "Unexpected" }
+
+ !$acc parallel default (private) ! { dg-error "Unclassifiable OpenACC directive" }
+ !$acc end parallel ! { dg-error "Unexpected" }
+
+ !$acc parallel default (none)
+ !$acc end parallel
+
+ !$acc parallel default (firstprivate) ! { dg-error "Unclassifiable OpenACC directive" }
+ !$acc end parallel ! { dg-error "Unexpected" }
+end program tile
new file mode 100644
@@ -0,0 +1,59 @@
+! Ensure that the internal array variables, offset, lbound, etc., don't
+! trigger errors with default(none).
+
+! { dg-do compile }
+
+program main
+ implicit none
+ integer i
+ integer,parameter :: n = 100
+ integer,allocatable :: a1(:), a2(:,:)
+
+ allocate (a1 (n))
+ allocate (a2 (-n:n,-n:n))
+ a1 (:) = -1
+
+ !$acc parallel loop default(none) copy (a1(1:n))
+ do i = 1,n
+ a1(i) = i
+ end do
+ !$acc end parallel loop
+
+ call foo (a1)
+ call bar (a1, n)
+ call foobar (a2,n)
+
+contains
+
+ subroutine foo (da1)
+ integer :: da1(n)
+
+ !$acc parallel loop default(none) copy (da1(1:n))
+ do i = 1,n
+ da1(i) = i*2
+ end do
+ !$acc end parallel loop
+ end subroutine foo
+end program main
+
+subroutine bar (da2,n)
+ integer :: n, da2(n)
+ integer i
+
+ !$acc parallel loop default(none) copy (da2(1:n)) firstprivate(n)
+ do i = 1,n
+ da2(i) = i*3
+ end do
+ !$acc end parallel loop
+end subroutine bar
+
+subroutine foobar (da3,n)
+ integer :: n, da3(-n:n,-n:n)
+ integer i
+
+ !$acc parallel loop default(none) copy (da3(-n:n,-n:n)) firstprivate(n)
+ do i = 1,n
+ da3(i,0) = i*3
+ end do
+ !$acc end parallel loop
+end subroutine foobar
new file mode 100644
@@ -0,0 +1,11 @@
+! { dg-do compile }
+
+program test
+ integer a, b(100)
+
+ !$acc parallel firstprivate (a, b)
+ !$acc end parallel
+
+ !$acc parallel firstprivate (b(10:20)) ! { dg-error "Syntax error in OpenMP variable list" }
+ !$acc end parallel ! { dg-error "Unexpected !\\\$ACC END PARALLEL statement" }
+end program test
new file mode 100644
@@ -0,0 +1,68 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-omplower" }
+
+program main
+ integer, parameter :: n = 100
+ integer i, a(n), b(n)
+
+ do i = 1, n
+ b(i) = i
+ end do
+
+ !$acc parallel loop gang (static:*) num_gangs (10)
+ do i = 1, n
+ a(i) = b(i) + 0
+ end do
+ !$acc end parallel loop
+
+ call test (a, b, 0, n)
+
+ !$acc parallel loop gang (static:1) num_gangs (10)
+ do i = 1, n
+ a(i) = b(i) + 1
+ end do
+ !$acc end parallel loop
+
+ call test (a, b, 1, n)
+
+ !$acc parallel loop gang (static:2) num_gangs (10)
+ do i = 1, n
+ a(i) = b(i) + 2
+ end do
+ !$acc end parallel loop
+
+ call test (a, b, 2, n)
+
+ !$acc parallel loop gang (static:5) num_gangs (10)
+ do i = 1, n
+ a(i) = b(i) + 5
+ end do
+ !$acc end parallel loop
+
+ call test (a, b, 5, n)
+
+ !$acc parallel loop gang (static:20) num_gangs (10)
+ do i = 1, n
+ a(i) = b(i) + 20
+ end do
+ !$acc end parallel loop
+
+ call test (a, b, 20, n)
+
+end program main
+
+subroutine test (a, b, sarg, n)
+ integer n
+ integer a (n), b(n), sarg
+ integer i
+
+ do i = 1, n
+ if (a(i) .ne. b(i) + sarg) call abort ()
+ end do
+end subroutine test
+
+! { dg-final { scan-tree-dump-times "gang\\(static:\\\*\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "gang\\(static:1\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "gang\\(static:2\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "gang\\(static:5\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "gang\\(static:20\\)" 1 "omplower" } }
new file mode 100644
@@ -0,0 +1,23 @@
+! { dg-additional-options "-O2" }
+! { dg-additional-options "-ftree-parallelize-loops=32" }
+
+program main
+ implicit none
+
+ integer :: a(100,100), b(100,100)
+ integer :: i, j, d
+
+ !$acc kernels
+ do i=1,100
+ do j=1,100
+ a(i,j) = 1
+ b(i,j) = 2
+ a(i,j) = a(i,j) + b(i,j)
+ end do
+ end do
+ !$acc end kernels
+
+ d = sum(a)
+
+ print *,d
+end program main
new file mode 100644
@@ -0,0 +1,19 @@
+! { dg-additional-options "-O2" }
+! { dg-additional-options "-ftree-parallelize-loops=10" }
+
+program main
+ implicit none
+
+ integer :: a(10000), b(10000)
+ integer :: d
+
+ !$acc kernels
+ a = 1
+ b = 2
+ a = a + b
+ !$acc end kernels
+
+ d = sum(a)
+
+ print *,d
+end program main
@@ -5,7 +5,7 @@ program test
implicit none
integer :: i, j, k, l, a(10)
- common /b/ j, k
+ common /b/ k
real, pointer :: p1 => NULL()
complex :: c, d(10)
@@ -108,4 +108,4 @@ program test
!$acc host_data use_device(p1) ! { dg-error "POINTER" }
!$acc end host_data
-end program test
\ No newline at end of file
+end program test
@@ -66,7 +66,7 @@ program test
!$acc loop seq worker ! { dg-error "conflicts with" }
DO i = 1,10
ENDDO
- !$acc loop gang worker ! { dg-error "conflicts with" }
+ !$acc loop gang worker
DO i = 1,10
ENDDO
@@ -94,10 +94,10 @@ program test
!$acc loop seq vector ! { dg-error "conflicts with" }
DO i = 1,10
ENDDO
- !$acc loop gang vector ! { dg-error "conflicts with" }
+ !$acc loop gang vector
DO i = 1,10
ENDDO
- !$acc loop worker vector ! { dg-error "conflicts with" }
+ !$acc loop worker vector
DO i = 1,10
ENDDO
@@ -239,7 +239,7 @@ program test
!$acc loop seq worker ! { dg-error "conflicts with" }
DO i = 1,10
ENDDO
- !$acc loop gang worker ! { dg-error "conflicts with" }
+ !$acc loop gang worker
DO i = 1,10
ENDDO
@@ -267,10 +267,10 @@ program test
!$acc loop seq vector ! { dg-error "conflicts with" }
DO i = 1,10
ENDDO
- !$acc loop gang vector ! { dg-error "conflicts with" }
+ !$acc loop gang vector
DO i = 1,10
ENDDO
- !$acc loop worker vector ! { dg-error "conflicts with" }
+ !$acc loop worker vector
DO i = 1,10
ENDDO
@@ -392,7 +392,7 @@ program test
!$acc kernels loop seq worker ! { dg-error "conflicts with" }
DO i = 1,10
ENDDO
- !$acc kernels loop gang worker ! { dg-error "conflicts with" }
+ !$acc kernels loop gang worker
DO i = 1,10
ENDDO
@@ -420,10 +420,10 @@ program test
!$acc kernels loop seq vector ! { dg-error "conflicts with" }
DO i = 1,10
ENDDO
- !$acc kernels loop gang vector ! { dg-error "conflicts with" }
+ !$acc kernels loop gang vector
DO i = 1,10
ENDDO
- !$acc kernels loop worker vector ! { dg-error "conflicts with" }
+ !$acc kernels loop worker vector
DO i = 1,10
ENDDO
@@ -544,7 +544,7 @@ program test
!$acc parallel loop seq worker ! { dg-error "conflicts with" }
DO i = 1,10
ENDDO
- !$acc parallel loop gang worker ! { dg-error "conflicts with" }
+ !$acc parallel loop gang worker
DO i = 1,10
ENDDO
@@ -572,10 +572,10 @@ program test
!$acc parallel loop seq vector ! { dg-error "conflicts with" }
DO i = 1,10
ENDDO
- !$acc parallel loop gang vector ! { dg-error "conflicts with" }
+ !$acc parallel loop gang vector
DO i = 1,10
ENDDO
- !$acc parallel loop worker vector ! { dg-error "conflicts with" }
+ !$acc parallel loop worker vector
DO i = 1,10
ENDDO
@@ -646,4 +646,4 @@ program test
!$acc parallel loop gang worker tile(*)
DO i = 1,10
ENDDO
-end
\ No newline at end of file
+end
new file mode 100644
@@ -0,0 +1,7 @@
+! Ensure that loops not affiliated with acc compute regions cause an error.
+
+subroutine test1
+ !$acc loop gang ! { dg-error "loop directive must be associated with an OpenACC compute region" }
+ DO i = 1,10
+ ENDDO
+end subroutine test1
new file mode 100644
@@ -0,0 +1,363 @@
+! { dg-do compile }
+! { dg-additional-options "-fmax-errors=100" }
+
+! { dg-prune-output "sorry, unimplemented" }
+! { dg-prune-output "Error: work-sharing region" }
+
+program test
+ implicit none
+ integer :: i, j
+
+ !$acc kernels
+ !$acc loop auto
+ DO i = 1,10
+ ENDDO
+ !$acc loop gang
+ DO i = 1,10
+ ENDDO
+ !$acc loop gang(5)
+ DO i = 1,10
+ ENDDO
+ !$acc loop gang(num:5)
+ DO i = 1,10
+ ENDDO
+ !$acc loop gang(static:5)
+ DO i = 1,10
+ ENDDO
+ !$acc loop gang(static:*)
+ DO i = 1,10
+ ENDDO
+ !$acc loop gang
+ DO i = 1,10
+ !$acc loop vector
+ DO j = 1,10
+ ENDDO
+ !$acc loop worker
+ DO j = 1,10
+ ENDDO
+ ENDDO
+
+ !$acc loop worker
+ DO i = 1,10
+ ENDDO
+ !$acc loop worker(5)
+ DO i = 1,10
+ ENDDO
+ !$acc loop worker(num:5)
+ DO i = 1,10
+ ENDDO
+ !$acc loop worker
+ DO i = 1,10
+ !$acc loop vector
+ DO j = 1,10
+ ENDDO
+ ENDDO
+ !$acc loop gang worker
+ DO i = 1,10
+ ENDDO
+
+ !$acc loop vector
+ DO i = 1,10
+ ENDDO
+ !$acc loop vector(5)
+ DO i = 1,10
+ ENDDO
+ !$acc loop vector(length:5)
+ DO i = 1,10
+ ENDDO
+ !$acc loop vector
+ DO i = 1,10
+ ENDDO
+ !$acc loop gang vector
+ DO i = 1,10
+ ENDDO
+ !$acc loop worker vector
+ DO i = 1,10
+ ENDDO
+
+ !$acc loop auto
+ DO i = 1,10
+ ENDDO
+
+ !$acc loop tile(1)
+ DO i = 1,10
+ ENDDO
+ !$acc loop tile(2)
+ DO i = 1,10
+ ENDDO
+ !$acc loop tile(6-2)
+ DO i = 1,10
+ ENDDO
+ !$acc loop tile(6+2)
+ DO i = 1,10
+ ENDDO
+ !$acc loop tile(*)
+ DO i = 1,10
+ ENDDO
+ !$acc loop tile(*, 1)
+ DO i = 1,10
+ DO j = 1,10
+ ENDDO
+ ENDDO
+ !$acc loop tile(-1) ! { dg-warning "must be positive" }
+ do i = 1,10
+ enddo
+ !$acc loop vector tile(*)
+ DO i = 1,10
+ ENDDO
+ !$acc loop worker tile(*)
+ DO i = 1,10
+ ENDDO
+ !$acc loop gang tile(*)
+ DO i = 1,10
+ ENDDO
+ !$acc loop vector gang tile(*)
+ DO i = 1,10
+ ENDDO
+ !$acc loop vector worker tile(*)
+ DO i = 1,10
+ ENDDO
+ !$acc loop gang worker tile(*)
+ DO i = 1,10
+ ENDDO
+ !$acc end kernels
+
+
+ !$acc parallel
+ !$acc loop tile(1)
+ DO i = 1,10
+ ENDDO
+ !$acc loop tile(*)
+ DO i = 1,10
+ ENDDO
+ !$acc loop tile(2)
+ DO i = 1,10
+ DO j = 1,10
+ ENDDO
+ ENDDO
+ !$acc loop tile(-1) ! { dg-warning "must be positive" }
+ do i = 1,10
+ enddo
+ !$acc loop vector tile(*)
+ DO i = 1,10
+ ENDDO
+ !$acc loop worker tile(*)
+ DO i = 1,10
+ ENDDO
+ !$acc loop gang tile(*)
+ DO i = 1,10
+ ENDDO
+ !$acc loop vector gang tile(*)
+ DO i = 1,10
+ ENDDO
+ !$acc loop vector worker tile(*)
+ DO i = 1,10
+ ENDDO
+ !$acc loop gang worker tile(*)
+ DO i = 1,10
+ ENDDO
+ !$acc end parallel
+
+ !$acc kernels loop auto
+ DO i = 1,10
+ ENDDO
+ !$acc kernels loop gang
+ DO i = 1,10
+ ENDDO
+ !$acc kernels loop gang(5)
+ DO i = 1,10
+ ENDDO
+ !$acc kernels loop gang(num:5)
+ DO i = 1,10
+ ENDDO
+ !$acc kernels loop gang(static:5)
+ DO i = 1,10
+ ENDDO
+ !$acc kernels loop gang(static:*)
+ DO i = 1,10
+ ENDDO
+ !$acc kernels loop gang
+ DO i = 1,10
+ !$acc kernels loop gang ! { dg-error "OpenACC construct inside of non-OpenACC region" }
+ DO j = 1,10
+ ENDDO
+ ENDDO
+
+ !$acc kernels loop worker
+ DO i = 1,10
+ ENDDO
+ !$acc kernels loop worker(5)
+ DO i = 1,10
+ ENDDO
+ !$acc kernels loop worker(num:5)
+ DO i = 1,10
+ ENDDO
+ !$acc kernels loop worker
+ DO i = 1,10
+ !$acc kernels loop worker ! { dg-error "OpenACC construct inside of non-OpenACC region" }
+ DO j = 1,10
+ ENDDO
+ !$acc kernels loop gang ! { dg-error "OpenACC construct inside of non-OpenACC region" }
+ DO j = 1,10
+ ENDDO
+ ENDDO
+ !$acc kernels loop gang worker
+ DO i = 1,10
+ ENDDO
+
+ !$acc kernels loop vector
+ DO i = 1,10
+ ENDDO
+ !$acc kernels loop vector(5)
+ DO i = 1,10
+ ENDDO
+ !$acc kernels loop vector(length:5)
+ DO i = 1,10
+ ENDDO
+ !$acc kernels loop vector
+ DO i = 1,10
+ !$acc kernels loop vector ! { dg-error "OpenACC construct inside of non-OpenACC region" }
+ DO j = 1,10
+ ENDDO
+ !$acc kernels loop worker ! { dg-error "OpenACC construct inside of non-OpenACC region" }
+ DO j = 1,10
+ ENDDO
+ !$acc kernels loop gang ! { dg-error "OpenACC construct inside of non-OpenACC region" }
+ DO j = 1,10
+ ENDDO
+ ENDDO
+ !$acc kernels loop gang vector
+ DO i = 1,10
+ ENDDO
+ !$acc kernels loop worker vector
+ DO i = 1,10
+ ENDDO
+
+ !$acc kernels loop auto
+ DO i = 1,10
+ ENDDO
+
+ !$acc kernels loop tile(1)
+ DO i = 1,10
+ ENDDO
+ !$acc kernels loop tile(*)
+ DO i = 1,10
+ ENDDO
+ !$acc kernels loop tile(*, 1)
+ DO i = 1,10
+ DO j = 1,10
+ ENDDO
+ ENDDO
+ !$acc kernels loop tile(-1) ! { dg-warning "must be positive" }
+ do i = 1,10
+ enddo
+ !$acc kernels loop vector tile(*)
+ DO i = 1,10
+ ENDDO
+ !$acc kernels loop worker tile(*)
+ DO i = 1,10
+ ENDDO
+ !$acc kernels loop gang tile(*)
+ DO i = 1,10
+ ENDDO
+ !$acc kernels loop vector gang tile(*)
+ DO i = 1,10
+ ENDDO
+ !$acc kernels loop vector worker tile(*)
+ DO i = 1,10
+ ENDDO
+ !$acc kernels loop gang worker tile(*)
+ DO i = 1,10
+ ENDDO
+
+ !$acc parallel loop auto
+ DO i = 1,10
+ ENDDO
+ !$acc parallel loop gang
+ DO i = 1,10
+ ENDDO
+ !$acc parallel loop gang(static:5)
+ DO i = 1,10
+ ENDDO
+ !$acc parallel loop gang(static:*)
+ DO i = 1,10
+ ENDDO
+ !$acc parallel loop gang
+ DO i = 1,10
+ !$acc parallel loop gang ! { dg-error "OpenACC construct inside of non-OpenACC region" }
+ DO j = 1,10
+ ENDDO
+ ENDDO
+
+ !$acc parallel loop worker
+ DO i = 1,10
+ ENDDO
+ !$acc parallel loop worker
+ DO i = 1,10
+ !$acc parallel loop worker ! { dg-error "OpenACC construct inside of non-OpenACC region" }
+ DO j = 1,10
+ ENDDO
+ !$acc parallel loop gang ! { dg-error "OpenACC construct inside of non-OpenACC region" }
+ DO j = 1,10
+ ENDDO
+ ENDDO
+ !$acc parallel loop gang worker
+ DO i = 1,10
+ ENDDO
+
+ !$acc parallel loop vector
+ DO i = 1,10
+ !$acc parallel loop vector ! { dg-error "OpenACC construct inside of non-OpenACC region" }
+ DO j = 1,10
+ ENDDO
+ !$acc parallel loop worker ! { dg-error "OpenACC construct inside of non-OpenACC region" }
+ DO j = 1,10
+ ENDDO
+ !$acc parallel loop gang ! { dg-error "OpenACC construct inside of non-OpenACC region" }
+ DO j = 1,10
+ ENDDO
+ ENDDO
+ !$acc parallel loop gang vector
+ DO i = 1,10
+ ENDDO
+ !$acc parallel loop worker vector
+ DO i = 1,10
+ ENDDO
+
+ !$acc parallel loop auto
+ DO i = 1,10
+ ENDDO
+
+ !$acc parallel loop tile(1)
+ DO i = 1,10
+ ENDDO
+ !$acc parallel loop tile(*)
+ DO i = 1,10
+ ENDDO
+ !$acc parallel loop tile(*, 1)
+ DO i = 1,10
+ DO j = 1,10
+ ENDDO
+ ENDDO
+ !$acc parallel loop tile(-1) ! { dg-warning "must be positive" }
+ do i = 1,10
+ enddo
+ !$acc parallel loop vector tile(*)
+ DO i = 1,10
+ ENDDO
+ !$acc parallel loop worker tile(*)
+ DO i = 1,10
+ ENDDO
+ !$acc parallel loop gang tile(*)
+ DO i = 1,10
+ ENDDO
+ !$acc parallel loop vector gang tile(*)
+ DO i = 1,10
+ ENDDO
+ !$acc parallel loop vector worker tile(*)
+ DO i = 1,10
+ ENDDO
+ !$acc parallel loop gang worker tile(*)
+ DO i = 1,10
+ ENDDO
+end
new file mode 100644
@@ -0,0 +1,82 @@
+! { dg-do compile }
+! { dg-additional-options "-fmax-errors=100" }
+
+! This error is temporary. Remove when support is added for these clauses
+! in the middle end.
+! { dg-prune-output "sorry, unimplemented" }
+! { dg-prune-output "Error: work-sharing region" }
+
+program test
+ implicit none
+ integer :: i, j
+
+ !$acc parallel
+ !$acc loop auto
+ DO i = 1,10
+ ENDDO
+ !$acc loop gang
+ DO i = 1,10
+ ENDDO
+ !$acc loop gang(static:5)
+ DO i = 1,10
+ ENDDO
+ !$acc loop gang(static:*)
+ DO i = 1,10
+ ENDDO
+ !$acc loop gang
+ DO i = 1,10
+ !$acc loop vector
+ DO j = 1,10
+ ENDDO
+ !$acc loop worker
+ DO j = 1,10
+ ENDDO
+ ENDDO
+
+ !$acc loop worker
+ DO i = 1,10
+ ENDDO
+ !$acc loop worker
+ DO i = 1,10
+ !$acc loop vector
+ DO j = 1,10
+ ENDDO
+ ENDDO
+ !$acc loop gang worker
+ DO i = 1,10
+ ENDDO
+
+ !$acc loop vector
+ DO i = 1,10
+ ENDDO
+ !$acc loop vector(5) ! { dg-error "argument not permitted" }
+ DO i = 1,10
+ ENDDO
+ !$acc loop vector(length:5) ! { dg-error "argument not permitted" }
+ DO i = 1,10
+ ENDDO
+ !$acc loop vector
+ DO i = 1,10
+ ENDDO
+ !$acc loop gang vector
+ DO i = 1,10
+ ENDDO
+ !$acc loop worker vector
+ DO i = 1,10
+ ENDDO
+
+ !$acc loop auto
+ DO i = 1,10
+ ENDDO
+ !$acc end parallel
+
+ !$acc parallel loop vector
+ DO i = 1,10
+ ENDDO
+ !$acc parallel loop vector(5) ! { dg-error "argument not permitted" }
+ DO i = 1,10
+ ENDDO
+ !$acc parallel loop vector(length:5) ! { dg-error "argument not permitted" }
+ DO i = 1,10
+ ENDDO
+end
@@ -3,6 +3,11 @@
! test for tree-dump-original and spaces-commas
+! This error is temporary. Remove when support is added for these clauses
+! in the middle end.
+! { dg-prune-output "sorry, unimplemented" }
+! { dg-prune-output "Error: work-sharing region" }
+
program test
implicit none
integer :: i, j, k, m, sum
@@ -17,7 +22,7 @@ program test
!$acc loop independent gang (3)
DO i = 1,10
- !$acc loop worker(3) ! { dg-error "work-sharing region may not be closely nested inside of work-sharing, critical, ordered, master or explicit task region" }
+ !$acc loop worker(3)
DO j = 1,10
!$acc loop vector(5)
DO k = 1,10
new file mode 100644
@@ -0,0 +1,13 @@
+! Test if variable appearing in multiple clauses are errors.
+
+! { dg-compile }
+
+program combined
+ implicit none
+ integer a(100), i, j
+
+ !$acc parallel loop reduction (+:j) copy (j) copyout(j) ! { dg-error "Symbol 'j' present on multiple clauses" }
+ do i = 1, 100
+ end do
+ !$acc end parallel loop
+end program combined
@@ -37,4 +37,3 @@ end program test
! { dg-final { scan-tree-dump-times "map\\(force_deviceptr:u\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "private\\(v\\)" 1 "original" } }
-! { dg-final { scan-tree-dump-times "firstprivate\\(w\\)" 1 "original" } }
new file mode 100644
@@ -0,0 +1,5 @@
+! { dg-do compile }
+
+program foo
+ !$acc update ! { dg-error "must contain at least one 'device' or 'host' or 'self' clause" }
+end program foo
@@ -432,7 +432,10 @@ enum omp_clause_code {
OMP_CLAUSE_NUM_WORKERS,
/* OpenACC clause: vector_length (integer-expression). */
- OMP_CLAUSE_VECTOR_LENGTH
+ OMP_CLAUSE_VECTOR_LENGTH,
+
+ /* OpenACC clause: tile ( size-expr-list ). */
+ OMP_CLAUSE_TILE
};
#undef DEFTREESTRUCT
@@ -932,6 +932,12 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, int flags)
case OMP_CLAUSE_INDEPENDENT:
pp_string (pp, "independent");
break;
+ case OMP_CLAUSE_TILE:
+ pp_string (pp, "tile(");
+ dump_generic_node (pp, OMP_CLAUSE_TILE_LIST (clause),
+ spc, flags, false);
+ pp_right_paren (pp);
+ break;
default:
/* Should never happen. */
@@ -328,6 +328,7 @@ unsigned const char omp_clause_num_ops[] =
1, /* OMP_CLAUSE_NUM_GANGS */
1, /* OMP_CLAUSE_NUM_WORKERS */
1, /* OMP_CLAUSE_VECTOR_LENGTH */
+ 1, /* OMP_CLAUSE_TILE */
};
const char * const omp_clause_code_name[] =
@@ -398,7 +399,8 @@ const char * const omp_clause_code_name[] =
"vector",
"num_gangs",
"num_workers",
- "vector_length"
+ "vector_length",
+ "tile"
};
@@ -11595,6 +11597,7 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data,
case OMP_CLAUSE_DEFAULTMAP:
case OMP_CLAUSE_AUTO:
case OMP_CLAUSE_SEQ:
+ case OMP_CLAUSE_TILE:
WALK_SUBTREE_TAIL (OMP_CLAUSE_CHAIN (*tp));
case OMP_CLAUSE_LASTPRIVATE:
@@ -1556,6 +1556,9 @@ extern void protected_set_expr_location (tree, location_t);
#define OMP_CLAUSE_DEFAULT_KIND(NODE) \
(OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_DEFAULT)->omp_clause.subcode.default_kind)
+#define OMP_CLAUSE_TILE_LIST(NODE) \
+ OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_TILE), 0)
+
/* SSA_NAME accessors. */
/* Returns the IDENTIFIER_NODE giving the SSA name a name or NULL_TREE