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 <ce...@codesourcery.com> Thomas Schwinge <tho...@codesourcery.com> James Norris <jnor...@codesourcery.com> Joseph Myers <jos...@codesourcery.com> Julian Brown <jul...@codesourcery.com> Bernd Schmidt <bschm...@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 <ce...@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; +}