2015-11-04 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 and independent as loop clauses. Associate REDUCTION
clauses with parallel and kernels.
* c-pragma.h (enum pragma_omp_clause): Add entries for
PRAGMA_OACC_CLAUSE_{INDEPENDENT,TILE,DEFAULT}.
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-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}.
* semantics.c (finish_omp_clauses): Add support for
OMP_CLAUSE_INDEPENDENT and OMP_CLAUSE_TILE.
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.
@@ -709,12 +709,20 @@ 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:
OMP_CLAUSE_CHAIN (clauses) = loop_clauses;
loop_clauses = clauses;
break;
+ /* 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,
@@ -10006,6 +10006,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;
@@ -10102,6 +10104,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;
@@ -10539,10 +10543,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;
@@ -10563,7 +10570,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;
@@ -10577,7 +10584,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 %<)%>");
@@ -11234,7 +11244,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;
@@ -11376,6 +11389,83 @@ 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;
+
+ vec<tree, va_gc> *tvec = make_tree_vector ();
+
+ 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)
+ goto cleanup_error;
+
+ 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;
+ }
+ }
+
+ vec_safe_push (tvec, expr);
+ 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 = build_tree_list_vec (tvec);
+ OMP_CLAUSE_TILE_LIST (c) = tile;
+ OMP_CLAUSE_CHAIN (c) = list;
+ release_tree_vector (tvec);
+ return c;
+
+ cleanup_error:
+ c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, "expected %<)%>");
+ release_tree_vector (tvec);
+ return list;
+}
+
/* OpenACC:
wait ( int-expr-list ) */
@@ -12576,6 +12666,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";
@@ -12601,6 +12695,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";
@@ -12646,6 +12745,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,
@@ -12727,7 +12830,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:
@@ -13140,13 +13243,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)
@@ -13191,6 +13296,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) \
@@ -13206,8 +13312,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) \
@@ -12946,10 +12946,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;
@@ -29119,6 +29119,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;
@@ -29213,6 +29215,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;
@@ -29645,7 +29649,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;
@@ -29713,6 +29720,59 @@ 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;
+
+ vec<tree, va_gc> *tvec = make_tree_vector ();
+
+ 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)
+ goto cleanup_error;
+
+ vec_safe_push (tvec, expr);
+
+ 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 = build_tree_list_vec (tvec);
+ OMP_CLAUSE_TILE_LIST (c) = tile;
+ OMP_CLAUSE_CHAIN (c) = list;
+ release_tree_vector (tvec);
+ return c;
+
+ cleanup_error:
+ release_tree_vector (tvec);
+ return list;
+}
+
/* OpenACC:
vector_length ( expression ) */
@@ -29859,10 +29919,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 +29947,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 +29961,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))
@@ -31466,6 +31533,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";
@@ -31474,6 +31545,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,
@@ -31487,6 +31563,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:
clauses = cp_parser_omp_clause_num_gangs (parser, clauses);
c_name = "num_gangs";
@@ -31515,6 +31597,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";
@@ -31528,6 +31615,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,
@@ -31615,7 +31706,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:
@@ -34459,12 +34550,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,
@@ -34509,6 +34603,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) \
@@ -34524,7 +34619,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) \
@@ -34533,6 +34630,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) )
@@ -6704,9 +6704,51 @@ 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:
+ {
+ tree list = OMP_CLAUSE_TILE_LIST (c);
+
+ while (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;
+ list = TREE_CHAIN (list);
+ }
+ }
+ break;
+
case OMP_CLAUSE_INBRANCH:
case OMP_CLAUSE_NOTINBRANCH:
if (branch_seen)
@@ -6995,9 +6995,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;
@@ -7007,6 +7016,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:
@@ -7482,6 +7492,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:
@@ -2230,6 +2230,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:
@@ -2408,6 +2409,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,117 @@
+// { dg-do compile }
+// { dg-options "-fopenacc -fdump-tree-gimple" }
+
+// { 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" } }
+// { dg-final { scan-tree-dump-times "reduction..:z. map.force_tofrom:z .len: 4.." 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;
+}
@@ -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
@@ -905,6 +905,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:
@@ -1561,6 +1561,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