On 10/23/2015 02:31 PM, Cesar Philippidis wrote: > On 10/23/2015 01:31 PM, Jakub Jelinek wrote: >> On Fri, Oct 23, 2015 at 01:17:07PM -0700, Cesar Philippidis wrote: >>> Good idea, thanks. This patch also corrects the problems parsing weird >>> combinations of num, static and length arguments that you mentioned >>> elsewhere. >>> >>> Is this OK for trunk? >> >> I'd strongly prefer to see always patches accompanied by testcases. >> >>> + loc = c_parser_peek_token (parser)->location; >>> + op_to_parse = &op0; >>> + >>> + if ((c_parser_next_token_is (parser, CPP_NAME) >>> + || c_parser_next_token_is (parser, CPP_KEYWORD)) >>> + && c_parser_peek_2nd_token (parser)->type == CPP_COLON) >>> + { >>> + tree name_kind = c_parser_peek_token (parser)->value; >>> + const char *p = IDENTIFIER_POINTER (name_kind); >> >> I think I'd prefer not to peek at this at all if it is RID_STATIC, >> so perhaps just have (and name_kind is weird): >> else >> { >> tree val = c_parser_peek_token (parser)->value; >> if (strcmp (id, IDENTIFIER_POINTER (val)) == 0) >> { >> c_parser_consume_token (parser); /* id */ >> c_parser_consume_token (parser); /* ':' */ >> } >> else >> { >> ... >> } >> } >> ? > > My plan over here was try and catch any arguments with a colon. But that > fell threw because... > >>> + if (kind == OMP_CLAUSE_GANG >>> + && c_parser_next_token_is_keyword (parser, RID_STATIC)) >>> + { >>> + c_parser_consume_token (parser); /* static */ >>> + c_parser_consume_token (parser); /* ':' */ >>> + >>> + op_to_parse = &op1; >>> + if (c_parser_next_token_is (parser, CPP_MULT)) >>> + { >>> + c_parser_consume_token (parser); >>> + *op_to_parse = integer_minus_one_node; >>> + >>> + /* Consume a comma if present. */ >>> + if (c_parser_next_token_is (parser, CPP_COMMA)) >>> + c_parser_consume_token (parser); >> >> Doesn't this mean that you happily parse >> gang (static: * abc) >> or >> gang (static:*num:1) >> etc.? I'd say the comma should be non-optional (i.e. either accept >> CPP_COMMA, or CPP_CLOSE_PARENT, but nothing else) in that case (at least, >> when in OpenMP grammar something is *-list it is meant to be comma >> separated). > > I'm not handling commas properly. My next patch is going to handle the > static argument separately. > >>> + /* Consume a comma if present. */ >>> + if (c_parser_next_token_is (parser, CPP_COMMA)) >>> + c_parser_consume_token (parser); >> >> Similarly this means >> gang (num: 5 static: *) >> is accepted. If it is valid, then again it should have testsuite coverage. > > I'll include a test case for this with the next patch.
Here's the updated patch. Hopefully I addressed everything. Thank you for suggesting all of those test cases. Is this OK for trunk? Cesar
2015-10-23 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-23 Cesar Philippidis <ce...@codesourcery.com> gcc/testsuite/ * c-c++-common/goacc/loop-shape.c (int main): diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index c8c6a2d..7d2baa9 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -11188,6 +11188,156 @@ 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 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 + { + next = c_parser_peek_token (parser); + loc = next->location; + 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); + } + + /* First extract any length:, num: and static: arguments. */ + + /* Gang static argument. */ + if (c_parser_next_token_is (parser, CPP_KEYWORD) + && c_parser_next_token_is_keyword (parser, RID_STATIC)) + { + if (kind != OMP_CLAUSE_GANG) + { + c_parser_error (parser, "invalid %<static%> argument"); + goto cleanup_error; + } + + c_parser_consume_token (parser); + + if (!c_parser_require (parser, CPP_COLON, "expected %<:%>")) + goto cleanup_error; + + op_to_parse = &op1; + if (*op_to_parse != NULL) + { + 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); + *op_to_parse = integer_minus_one_node; + continue; + } + } + /* 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 (*op_to_parse != 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); + *op_to_parse = expr; + } + while (c_parser_next_token_is (parser, CPP_COMMA) + && !c_parser_next_token_is (parser, CPP_CLOSE_PAREN)); + + 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); + if (op0) + OMP_CLAUSE_OPERAND (c, 0) = op0; + if (op1) + OMP_CLAUSE_OPERAND (c, 1) = op1; + 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 +12543,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 +12584,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 +12637,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 +12655,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 +13190,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..1053361 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/loop-shape.c @@ -0,0 +1,218 @@ +/* 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 "unexpected argument" } */ + 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 "unexpected argument" } */ + 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++) + ; + + + return 0; +}