From 765dd3918b6b52d766a5034080515acbdda738f0 Mon Sep 17 00:00:00 2001 From: Cesar Philippidis Date: Tue, 27 Oct 2015 13:21:31 -0700 Subject: [PATCH] c-parser.c (c_parser_oacc_shape_clause): New. 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. gcc/testsuite/ * c-c++-common/goacc/loop-shape.c: New test. Co-Authored-By: Bernd Schmidt Co-Authored-By: James Norris Co-Authored-By: Joseph Myers Co-Authored-By: Julian Brown Co-Authored-By: Thomas Schwinge From-SVN: r229461 --- gcc/c/ChangeLog | 12 + gcc/c/c-parser.c | 191 +++++++++++ gcc/testsuite/ChangeLog | 4 + gcc/testsuite/c-c++-common/goacc/loop-shape.c | 322 ++++++++++++++++++ 4 files changed, 529 insertions(+) create mode 100644 gcc/testsuite/c-c++-common/goacc/loop-shape.c diff --git a/gcc/c/ChangeLog b/gcc/c/ChangeLog index da4c8016834..78597bec268 100644 --- a/gcc/c/ChangeLog +++ b/gcc/c/ChangeLog @@ -1,3 +1,15 @@ +2015-10-27 Cesar Philippidis + Thomas Schwinge + James Norris + Joseph Myers + Julian Brown + Bernd Schmidt + + * 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-27 Thomas Schwinge James Norris Cesar Philippidis diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index ec918c14ca8..05700769449 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -11186,6 +11186,167 @@ c_parser_omp_clause_num_workers (c_parser *parser, tree list) return 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 % arguments"); + 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; + } + + location_t expr_loc = c_parser_peek_token (parser)->location; + tree expr = c_parser_expr_no_commas (parser, NULL).value; + if (expr == error_mark_node) + goto cleanup_error; + + mark_exp_read (expr); + expr = c_fully_fold (expr, false, NULL); + + /* Attempt to statically determine when the number isn't a + positive integer. */ + + if (!INTEGRAL_TYPE_P (TREE_TYPE (expr))) + { + c_parser_error (parser, "expected integer expression"); + return list; + } + + tree c = fold_build2_loc (expr_loc, LE_EXPR, boolean_type_node, expr, + build_int_cst (TREE_TYPE (expr), 0)); + if (c == boolean_true_node) + { + warning_at (loc, 0, + "%<%s%> value must be positive", str); + expr = integer_one_node; + } + + 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); + + if (ops[1]) + OMP_CLAUSE_OPERAND (c, 1) = ops[1]; + + OMP_CLAUSE_OPERAND (c, 0) = ops[0]; + 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, 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 )] */ @@ -12392,6 +12553,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"; @@ -12428,6 +12594,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"; @@ -12476,6 +12647,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"; @@ -12484,6 +12665,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; @@ -12960,6 +13146,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/ChangeLog b/gcc/testsuite/ChangeLog index f8ef47d8b50..0eec83b596a 100644 --- a/gcc/testsuite/ChangeLog +++ b/gcc/testsuite/ChangeLog @@ -1,3 +1,7 @@ +2015-10-27 Cesar Philippidis + + * c-c++-common/goacc/loop-shape.c: New test. + 2015-10-27 Nathan Sidwell Cesar Philippidis 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 00000000000..b6d3156074d --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/loop-shape.c @@ -0,0 +1,322 @@ +/* Exercise *_parser_oacc_shape_clause by checking various combinations + of gang, worker and vector clause arguments. */ + +/* { dg-compile } */ + +int main () +{ + int i; + int v = 32, w = 19; + int length = 1, num = 5; + + /* 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++) + ; + + #pragma acc kernels + #pragma acc loop gang(-12) /* { dg-warning "" } */ + for (i = 0; i < 10; i++) + ; + + #pragma acc kernels + #pragma acc loop gang(-1.0) /* { dg-error "" } */ + for (i = 0; i < 10; i++) + ; + + #pragma acc kernels + #pragma acc loop gang(1.0) /* { dg-error "" } */ + for (i = 0; i < 10; i++) + ; + + #pragma acc kernels + #pragma acc loop gang(num:-1.0) /* { dg-error "" } */ + for (i = 0; i < 10; i++) + ; + + #pragma acc kernels + #pragma acc loop gang(num:1.0) /* { dg-error "" } */ + for (i = 0; i < 10; i++) + ; + + #pragma acc kernels + #pragma acc loop gang(static:-1.0) /* { dg-error "" } */ + for (i = 0; i < 10; i++) + ; + + #pragma acc kernels + #pragma acc loop gang(static:1.0) /* { dg-error "" } */ + for (i = 0; i < 10; i++) + ; + + #pragma acc kernels + #pragma acc loop worker(-1.0) /* { dg-error "" } */ + for (i = 0; i < 10; i++) + ; + + #pragma acc kernels + #pragma acc loop worker(1.0) /* { dg-error "" } */ + for (i = 0; i < 10; i++) + ; + + #pragma acc kernels + #pragma acc loop worker(num:-1.0) /* { dg-error "" } */ + for (i = 0; i < 10; i++) + ; + + #pragma acc kernels + #pragma acc loop worker(num:1.0) /* { dg-error "" } */ + for (i = 0; i < 10; i++) + ; + + #pragma acc kernels + #pragma acc loop vector(-1.0) /* { dg-error "" } */ + for (i = 0; i < 10; i++) + ; + + #pragma acc kernels + #pragma acc loop vector(1.0) /* { dg-error "" } */ + for (i = 0; i < 10; i++) + ; + + #pragma acc kernels + #pragma acc loop vector(length:-1.0) /* { dg-error "" } */ + for (i = 0; i < 10; i++) + ; + + #pragma acc kernels + #pragma acc loop vector(length:1.0) /* { dg-error "" } */ + for (i = 0; i < 10; i++) + ; + + return 0; +} -- 2.30.2