This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
Re: [OpenACC 4/11] C FE changes
- From: Cesar Philippidis <cesar_philippidis at mentor dot com>
- To: Jakub Jelinek <jakub at redhat dot com>
- Cc: Nathan Sidwell <nathan at acm dot org>, GCC Patches <gcc-patches at gcc dot gnu dot org>, Bernd Schmidt <bschmidt at redhat dot com>, Jason Merrill <jason at redhat dot com>, "Joseph S. Myers" <joseph at codesourcery dot com>
- Date: Sat, 24 Oct 2015 14:10:26 -0700
- Subject: Re: [OpenACC 4/11] C FE changes
- Authentication-results: sourceware.org; auth=none
- References: <20151022082249 dot GF478 at tucnak dot redhat dot com> <562A95C3 dot 2040100 at mentor dot com> <20151023203155 dot GJ478 at tucnak dot redhat dot com> <562AA73D dot 6010809 at mentor dot com> <562AED97 dot 5070803 at mentor dot com> <20151024080356 dot GL478 at tucnak dot redhat dot com>
On 10/24/2015 01:03 AM, Jakub Jelinek wrote:
> On Fri, Oct 23, 2015 at 07:31:51PM -0700, Cesar Philippidis wrote:
>
>> +static tree
>> +c_parser_oacc_shape_clause (c_parser *parser, omp_clause_code kind,
>> + const char *str, tree list)
>> +{
>> + const char *id = "num";
>> + tree op0 = NULL_TREE, op1 = NULL_TREE, c;
>> + location_t loc = c_parser_peek_token (parser)->location;
>> +
>> + if (kind == OMP_CLAUSE_VECTOR)
>> + id = "length";
>> +
>> + if (c_parser_next_token_is (parser, CPP_OPEN_PAREN))
>> + {
>> + tree *op_to_parse = &op0;
>> + c_token *next;
>> +
>> + c_parser_consume_token (parser);
>> +
>> + do
>> + {
>> + op_to_parse = &op0;
>> +
>> + /* Consume a comma if present. */
>> + if (c_parser_next_token_is (parser, CPP_COMMA))
>> + {
>> + if (op0 == NULL && op1 == NULL)
>> + {
>> + c_parser_error (parser, "unexpected argument");
>> + goto cleanup_error;
>> + }
>> +
>> + c_parser_consume_token (parser);
>> + }
>
> This means you parse
> gang (, static: *)
> vector (, 5)
> etc., even when you error on it afterwards with unexpected argument,
> it is still different diagnostics from other invalid tokens immediately
> after the opening (.
So you didn't like how the error messages are inconsistent? It was
catching those errors.
I've added those new test cases. Unfortunately, c and c++ report
different error messages, so I had the make dg-error generic to that
line containing those types of errors.
> Also, loc and next are wrong if there is a valid comma.
Yeah, I don't think it needs to be adjusted in the loop. c_parser_error
already knows where to report the error at anyway.
> So I'm really wondering why
> gang (static: *, num: 5)
> works, because next is the CPP_COMMA token, so while
> c_parser_next_token_is (parser, CPP_NAME) matches the actual name,
> what exactly next->value contains is unclear.
>
> I think it would be better to:
>
> tree ops[2] = { NULL_TREE, NULL_TREE };
>
> do
> {
> // Note, declare these here
> c_token *next = c_parser_peek_token (parser);
> location_t loc = next->location;
> // Just use ops[idx] instead of *op_to_parse etc., though if you strongly
> // prefer *op_to_parse, I won't object.
> int idx = 0;
> // Note it seems generally the C parser doesn't check for CPP_KEYWORD
> // before calling c_parser_next_token_is_keyword. And I'd just do it
> // for OMP_CLAUSE_GANG, which has it in the grammar.
> if (kind == OMP_CLAUSE_GANG
> && c_parser_next_token_is_keyword (parser, RID_STATIC))
> {
> // ...
> // Your current code, except that for
> if (c_parser_next_token_is (parser, CPP_MULT))
> {
> c_parser_consume_token (parser);
> if (c_parser_next_token_is (parser, CPP_COMMA))
> {
> c_parser_consume_token (parser);
> continue;
> }
> break;
> }
> }
> else if (... num: / length: )
> {
> // ...
> }
> // ...
> mark_exp_read (expr);
> ops[idx] = expr;
>
> if (kind == OMP_CLAUSE_GANG
> && c_parser_next_token_is (parser, CPP_COMMA))
> {
> c_parser_consume_token (parser);
> continue;
> }
> break;
> }
> while (1);
>
> if (!c_parser_require (parser, CPP_CLOSE_PAREN, "expected %<)%>"))
> goto cleanup_error;
>
> That way you don't parse something that is not in the grammar.
I did that. It turned out to be a little more compact than what I had
before. Is this OK for trunk?
Cesar
2015-10-24 Cesar Philippidis <cesar@codesourcery.com>
Thomas Schwinge <thomas@codesourcery.com>
James Norris <jnorris@codesourcery.com>
Joseph Myers <joseph@codesourcery.com>
Julian Brown <julian@codesourcery.com>
Bernd Schmidt <bschmidt@redhat.com>
gcc/c/
* c-parser.c (c_parser_oacc_shape_clause): New.
(c_parser_oacc_simple_clause): New.
(c_parser_oacc_all_clauses): Add auto, gang, seq, vector, worker.
(OACC_LOOP_CLAUSE_MASK): Add gang, worker, vector, auto, seq.
2015-10-24 Cesar Philippidis <cesar@codesourcery.com>
gcc/testsuite/
* c-c++-common/goacc/loop-shape.c: New test.
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index c8c6a2d..2ad3825 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -11188,6 +11188,144 @@ c_parser_omp_clause_num_workers (c_parser *parser, tree list)
}
/* OpenACC:
+
+ gang [( gang-arg-list )]
+ worker [( [num:] int-expr )]
+ vector [( [length:] int-expr )]
+
+ where gang-arg is one of:
+
+ [num:] int-expr
+ static: size-expr
+
+ and size-expr may be:
+
+ *
+ int-expr
+*/
+
+static tree
+c_parser_oacc_shape_clause (c_parser *parser, omp_clause_code kind,
+ const char *str, tree list)
+{
+ const char *id = "num";
+ tree ops[2] = { NULL_TREE, NULL_TREE }, c;
+ location_t loc = c_parser_peek_token (parser)->location;
+
+ if (kind == OMP_CLAUSE_VECTOR)
+ id = "length";
+
+ if (c_parser_next_token_is (parser, CPP_OPEN_PAREN))
+ {
+ c_parser_consume_token (parser);
+
+ do
+ {
+ c_token *next = c_parser_peek_token (parser);
+ int idx = 0;
+
+ /* Gang static argument. */
+ if (kind == OMP_CLAUSE_GANG
+ && c_parser_next_token_is_keyword (parser, RID_STATIC))
+ {
+ c_parser_consume_token (parser);
+
+ if (!c_parser_require (parser, CPP_COLON, "expected %<:%>"))
+ goto cleanup_error;
+
+ idx = 1;
+ if (ops[idx] != NULL_TREE )
+ {
+ c_parser_error (parser, "too many %<static%> arguements");
+ goto cleanup_error;
+ }
+
+ /* Check for the '*' argument. */
+ if (c_parser_next_token_is (parser, CPP_MULT))
+ {
+ c_parser_consume_token (parser);
+ ops[idx] = integer_minus_one_node;
+
+ if (c_parser_next_token_is (parser, CPP_COMMA))
+ {
+ c_parser_consume_token (parser);
+ continue;
+ }
+ else
+ break;
+ }
+ }
+ /* Worker num: argument and vector length: arguments. */
+ else if (c_parser_next_token_is (parser, CPP_NAME)
+ && strcmp (id, IDENTIFIER_POINTER (next->value)) == 0
+ && c_parser_peek_2nd_token (parser)->type == CPP_COLON)
+ {
+ c_parser_consume_token (parser); /* id */
+ c_parser_consume_token (parser); /* ':' */
+ }
+
+ /* Now collect the actual argument. */
+ if (ops[idx] != NULL_TREE)
+ {
+ c_parser_error (parser, "unexpected argument");
+ goto cleanup_error;
+ }
+
+ tree expr = c_parser_expr_no_commas (parser, NULL).value;
+ if (expr == error_mark_node)
+ goto cleanup_error;
+
+ mark_exp_read (expr);
+ ops[idx] = expr;
+
+ if (kind == OMP_CLAUSE_GANG
+ && c_parser_next_token_is (parser, CPP_COMMA))
+ {
+ c_parser_consume_token (parser);
+ continue;
+ }
+ break;
+ }
+ while (1);
+
+ if (!c_parser_require (parser, CPP_CLOSE_PAREN, "expected %<)%>"))
+ goto cleanup_error;
+ }
+
+ check_no_duplicate_clause (list, kind, str);
+
+ c = build_omp_clause (loc, kind);
+ OMP_CLAUSE_OPERAND (c, 0) = ops[0];
+ if (ops[1])
+ OMP_CLAUSE_OPERAND (c, 1) = ops[1];
+ OMP_CLAUSE_CHAIN (c) = list;
+
+ return c;
+
+ cleanup_error:
+ c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, 0);
+ return list;
+}
+
+/* OpenACC:
+ auto
+ independent
+ nohost
+ seq */
+
+static tree
+c_parser_oacc_simple_clause (c_parser *parser ATTRIBUTE_UNUSED,
+ enum omp_clause_code code, tree list)
+{
+ check_no_duplicate_clause (list, code, omp_clause_code_name[code]);
+
+ tree c = build_omp_clause (c_parser_peek_token (parser)->location, code);
+ OMP_CLAUSE_CHAIN (c) = list;
+
+ return c;
+}
+
+/* OpenACC:
async [( int-expr )] */
static tree
@@ -12393,6 +12531,11 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
clauses = c_parser_oacc_clause_async (parser, clauses);
c_name = "async";
break;
+ case PRAGMA_OACC_CLAUSE_AUTO:
+ clauses = c_parser_oacc_simple_clause (parser, OMP_CLAUSE_AUTO,
+ clauses);
+ c_name = "auto";
+ break;
case PRAGMA_OACC_CLAUSE_COLLAPSE:
clauses = c_parser_omp_clause_collapse (parser, clauses);
c_name = "collapse";
@@ -12429,6 +12572,11 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
clauses = c_parser_omp_clause_firstprivate (parser, clauses);
c_name = "firstprivate";
break;
+ case PRAGMA_OACC_CLAUSE_GANG:
+ c_name = "gang";
+ clauses = c_parser_oacc_shape_clause (parser, OMP_CLAUSE_GANG,
+ c_name, clauses);
+ break;
case PRAGMA_OACC_CLAUSE_HOST:
clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "host";
@@ -12477,6 +12625,16 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "self";
break;
+ case PRAGMA_OACC_CLAUSE_SEQ:
+ clauses = c_parser_oacc_simple_clause (parser, OMP_CLAUSE_SEQ,
+ clauses);
+ c_name = "seq";
+ break;
+ case PRAGMA_OACC_CLAUSE_VECTOR:
+ c_name = "vector";
+ clauses = c_parser_oacc_shape_clause (parser, OMP_CLAUSE_VECTOR,
+ c_name, clauses);
+ break;
case PRAGMA_OACC_CLAUSE_VECTOR_LENGTH:
clauses = c_parser_omp_clause_vector_length (parser, clauses);
c_name = "vector_length";
@@ -12485,6 +12643,11 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
clauses = c_parser_oacc_clause_wait (parser, clauses);
c_name = "wait";
break;
+ case PRAGMA_OACC_CLAUSE_WORKER:
+ c_name = "worker";
+ clauses = c_parser_oacc_shape_clause (parser, OMP_CLAUSE_WORKER,
+ c_name, clauses);
+ break;
default:
c_parser_error (parser, "expected %<#pragma acc%> clause");
goto saw_error;
@@ -13015,6 +13178,11 @@ 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_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_SEQ) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) )
static tree
diff --git a/gcc/testsuite/c-c++-common/goacc/loop-shape.c b/gcc/testsuite/c-c++-common/goacc/loop-shape.c
new file mode 100644
index 0000000..aa75ac9
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/loop-shape.c
@@ -0,0 +1,247 @@
+/* Exercise *_parser_oacc_shape_clause by checking various combinations
+ of gang, worker and vector clause arguments. */
+
+/* { dg-compile } */
+
+int main ()
+{
+ int i;
+ int v, w;
+ int length, num;
+
+ /* Valid uses. */
+
+ #pragma acc kernels
+ #pragma acc loop gang worker vector
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(26)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(v)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop vector(length: 16)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop vector(length: v)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop vector(16)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop vector(v)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop worker(num: 16)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop worker(num: v)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop worker(16)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop worker(v)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(static: 16, num: 5)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(static: v, num: w)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop vector(length)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop worker(num)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(num, static: 6)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(static: 5, num)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(1, static:*)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(static:*, 1)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(1, static:*)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(num: 5, static: 4)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(num: v, static: w)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(num, static:num)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop vector(length:length)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop worker(num:length)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop worker(num:num)
+ for (i = 0; i < 10; i++)
+ ;
+
+ /* Invalid uses. */
+
+ #pragma acc kernels
+#pragma acc loop gang(16, 24) /* { dg-error "unexpected argument" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(v, w) /* { dg-error "unexpected argument" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(num: 1, num:2, num:3, 4) /* { dg-error "unexpected argument" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(num: 1 num:2, num:3, 4) /* { dg-error "expected '.' before" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(1, num:2, num:3, 4) /* { dg-error "unexpected argument" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(num, num:5) /* { dg-error "unexpected argument" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(length:num) /* { dg-error "" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop vector(5, length:length) /* { dg-error "expected '.' before" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop vector(num:length) /* { dg-error "" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop worker(length:5) /* { dg-error "expected '.' before" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop worker(1, num:2) /* { dg-error "expected '.' before" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(static: * abc) /* { dg-error "expected '.' before" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(static:*num:1) /* { dg-error "expected '.' before" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(num: 5 static: *) /* { dg-error "expected '.' before" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(,static: *) /* { dg-error "" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop vector(,length:5) /* { dg-error "" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop worker(,num:10) /* { dg-error "" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop worker(,10) /* { dg-error "" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop vector(,10) /* { dg-error "" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(,10) /* { dg-error "" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ return 0;
+}