From: Cesar Philippidis <cesar@codesourcery.com>
To: Jakub Jelinek <jakub@redhat.com>
Cc: "gcc-patches@gcc.gnu.org" <gcc-patches@gcc.gnu.org>
Subject: Re: [openacc] tile, independent, default, private and firstprivate support in c/++
Date: Wed, 04 Nov 2015 17:55:00 -0000 [thread overview]
Message-ID: <563A46A5.6000702@codesourcery.com> (raw)
In-Reply-To: <20151104102425.GB478@tucnak.redhat.com>
[-- Attachment #1: Type: text/plain, Size: 4123 bytes --]
On 11/04/2015 02:24 AM, Jakub Jelinek wrote:
> On Tue, Nov 03, 2015 at 02:16:59PM -0800, Cesar Philippidis wrote:
>> +
>> + do
>> + {
>> + if (c_parser_next_token_is (parser, CPP_MULT))
>> + {
>> + c_parser_consume_token (parser);
>> + expr = integer_minus_one_node;
>> + }
>> + else
>
> Is this right? If it is either * or (assignment) expression, then
> I'd expect to parse only CPP_MULT followed by CPP_CLOSE_PAREN
> or CPP_COMMA that way (C parser has 2 tokens look-ahead, so it should be
> fine), so that
> tile (a + b + c, *)
> is parsed as
> (a + b + c); -1
> and so is
> tile (*, a + b)
> as
> -1; (a + b)
> while
> tile (*a, *b)
> is
> *a; *b.
>
> Guess the gang clause parsing that went into the trunk already has the
> same bug,
> gang (static: *)
> or
> gang (static: *, num: 5)
> should be special, but
> gang (static: *ptr)
> should be
> gang (static: (*ptr))
Thanks for catching that. I'll fix the tile and shape scanners in both
front ends.
>> diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
>> index c73dcd0..14d006b 100644
>> --- a/gcc/cp/semantics.c
>> +++ b/gcc/cp/semantics.c
>> @@ -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;
>
> Have you verified pt.c does the right thing when instantiating the
> OMP_CLAUSE_TILE clause (I mean primarily the TREE_VEC in there)?
> There really should be testcases for that.
I don't think we have any support for templates in trunk yet. Should I
add it to this patch, or should I address that in a follow up patch?
By the way, template support for num_gangs, num_worker and vector_length
depend on this patch
<https://gcc.gnu.org/ml/gcc-patches/2015-10/msg03455.html> because the
c++ front end is currently incorrectly trying to do type checking as it
scans those clauses.
>> diff --git a/gcc/gimplify.c b/gcc/gimplify.c
>> index 03203c0..08b192d 100644
>> --- a/gcc/gimplify.c
>> +++ b/gcc/gimplify.c
>> @@ -6997,7 +6997,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
>>
>> case OMP_CLAUSE_DEVICE_RESIDENT:
>> case OMP_CLAUSE_USE_DEVICE:
>> - case OMP_CLAUSE_INDEPENDENT:
>> remove = true;
>> break;
>>
>> @@ -7007,6 +7006,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:
>> @@ -7014,6 +7014,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
>> case OMP_CLAUSE_NOGROUP:
>> case OMP_CLAUSE_THREADS:
>> case OMP_CLAUSE_SIMD:
>> + case OMP_CLAUSE_TILE:
>> break;
>
> No gimplification of the expressions in the tile clause?
Ultimately we're just ignoring it, but for completeness' sake I fixed this.
I'm still testing this patch, but something like this OK for trunk?
cesar
[-- Attachment #2: misc-c-cxx-clauses-4.diff --]
[-- Type: text/x-patch, Size: 36374 bytes --]
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.
diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c
index ca64eda..1306e09 100644
--- a/gcc/c-family/c-omp.c
+++ b/gcc/c-family/c-omp.c
@@ -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;
diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h
index 69e7392..953c4e3 100644
--- a/gcc/c-family/c-pragma.h
+++ b/gcc/c-family/c-pragma.h
@@ -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,
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index ec88c65..8360c35 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -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) \
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index 2363b9b..24cedf3 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -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;
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 12452e6..8f02f4c 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -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) )
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index c73dcd0..14d006b 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -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)
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 03203c0..91fde9b 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -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:
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index ccf0b63..9c87ac9 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -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:
diff --git a/gcc/testsuite/c-c++-common/goacc/combined-directives.c b/gcc/testsuite/c-c++-common/goacc/combined-directives.c
new file mode 100644
index 0000000..cefe712
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/combined-directives.c
@@ -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" } }
diff --git a/gcc/testsuite/c-c++-common/goacc/loop-clauses.c b/gcc/testsuite/c-c++-common/goacc/loop-clauses.c
new file mode 100644
index 0000000..97b8786
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/loop-clauses.c
@@ -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;
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/loop-shape.c b/gcc/testsuite/c-c++-common/goacc/loop-shape.c
index b6d3156..9708f7b 100644
--- a/gcc/testsuite/c-c++-common/goacc/loop-shape.c
+++ b/gcc/testsuite/c-c++-common/goacc/loop-shape.c
@@ -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++)
;
diff --git a/gcc/testsuite/c-c++-common/goacc/tile.c b/gcc/testsuite/c-c++-common/goacc/tile.c
new file mode 100644
index 0000000..2a81427
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/tile.c
@@ -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;
+}
diff --git a/gcc/tree-core.h b/gcc/tree-core.h
index 6b17da7..562f838 100644
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -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
diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index 8a2398d..6eeb12f 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -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. */
diff --git a/gcc/tree.c b/gcc/tree.c
index 18d6544..5b9a7bd 100644
--- a/gcc/tree.c
+++ b/gcc/tree.c
@@ -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:
diff --git a/gcc/tree.h b/gcc/tree.h
index 65c3117..be18b64a 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -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
next prev parent reply other threads:[~2015-11-04 17:55 UTC|newest]
Thread overview: 26+ messages / expand[flat|nested] mbox.gz Atom feed top
2015-11-03 22:17 Cesar Philippidis
2015-11-04 10:24 ` Jakub Jelinek
2015-11-04 17:55 ` Cesar Philippidis [this message]
2015-11-05 16:31 ` Jakub Jelinek
2015-11-05 4:58 ` Cesar Philippidis
2015-11-05 7:29 ` Jakub Jelinek
2015-11-05 14:58 ` Cesar Philippidis
2015-11-05 12:15 ` Thomas Schwinge
2015-11-05 14:48 ` Cesar Philippidis
2015-11-05 17:02 ` Thomas Schwinge
2015-11-05 17:13 ` Nathan Sidwell
2015-11-05 17:27 ` Cesar Philippidis
2015-11-06 2:11 ` Cesar Philippidis
2015-11-06 6:50 ` Jakub Jelinek
2015-11-06 13:44 ` Nathan Sidwell
2015-11-06 13:42 ` Nathan Sidwell
2015-11-09 11:31 ` [gomp4] " Thomas Schwinge
2021-07-21 22:06 ` Thomas Schwinge
[not found] ` <563D0735.7070601@mentor.com>
[not found] ` <87h9lqaofq.fsf@schwinge.name>
2015-11-06 23:31 ` [gomp4] backport trunk FE changes Cesar Philippidis
2015-11-07 11:45 ` Combined constructs' clause splitting (was: [gomp4] backport trunk FE changes) Thomas Schwinge
2015-11-07 16:14 ` Combined constructs' clause splitting Cesar Philippidis
2015-11-08 15:46 ` Tom de Vries
2015-11-19 0:03 ` Cesar Philippidis
2015-11-19 0:09 ` Tom de Vries
2015-11-07 12:30 ` [gomp4] backport trunk FE changes Thomas Schwinge
2015-11-07 16:05 ` Cesar Philippidis
Reply instructions:
You may reply publicly to this message via plain-text email
using any one of the following methods:
* Save the following mbox file, import it into your mail client,
and reply-to-all from there: mbox
Avoid top-posting and favor interleaved quoting:
https://en.wikipedia.org/wiki/Posting_style#Interleaved_style
* Reply using the --to, --cc, and --in-reply-to
switches of git-send-email(1):
git send-email \
--in-reply-to=563A46A5.6000702@codesourcery.com \
--to=cesar@codesourcery.com \
--cc=gcc-patches@gcc.gnu.org \
--cc=jakub@redhat.com \
/path/to/YOUR_REPLY
https://kernel.org/pub/software/scm/git/docs/git-send-email.html
* If your mail client supports setting the In-Reply-To header
via mailto: links, try the mailto: link
Be sure your reply has a Subject: header at the top and a blank line
before the message body.
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).