diff options
Diffstat (limited to 'gcc/c')
| -rw-r--r-- | gcc/c/ChangeLog | 23 | ||||
| -rw-r--r-- | gcc/c/c-parser.c | 123 | ||||
| -rw-r--r-- | gcc/c/c-typeck.c | 2 |
3 files changed, 139 insertions, 9 deletions
diff --git a/gcc/c/ChangeLog b/gcc/c/ChangeLog index 7839024..c734a81 100644 --- a/gcc/c/ChangeLog +++ b/gcc/c/ChangeLog @@ -1,3 +1,26 @@ +2015-11-05 Cesar Philippidis <cesar@codesourcery.com> + Thomas Schwinge <thomas@codesourcery.com> + James Norris <jnorris@codesourcery.com> + + * 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_parser_oacc_update): Update the error message for missing clauses. + * c-typeck.c (c_finish_omp_clauses): Add support for OMP_CLAUSE_TILE + and OMP_CLAUSE_INDEPENDENT. + 2015-11-05 Marek Polacek <polacek@redhat.com> PR c/68090 diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index e2f55b3..23d0107 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -10008,6 +10008,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; @@ -10104,6 +10106,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; @@ -10541,10 +10545,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; @@ -10565,7 +10572,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; @@ -10579,7 +10586,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 %<)%>"); @@ -11236,7 +11246,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; @@ -11378,6 +11391,79 @@ 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; + + 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) + { + c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, + "expected %<)%>"); + return list; + } + + 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; + } + } + + tile = tree_cons (NULL_TREE, expr, tile); + 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 = nreverse (tile); + OMP_CLAUSE_TILE_LIST (c) = tile; + OMP_CLAUSE_CHAIN (c) = list; + return c; +} + /* OpenACC: wait ( int-expr-list ) */ @@ -12605,6 +12691,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"; @@ -12630,6 +12720,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"; @@ -12675,6 +12770,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, @@ -12756,7 +12855,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: @@ -13169,13 +13268,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) @@ -13220,6 +13321,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) \ @@ -13235,8 +13337,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) \ @@ -13318,7 +13423,7 @@ c_parser_oacc_update (c_parser *parser) { error_at (loc, "%<#pragma acc update%> must contain at least one " - "%<device%> or %<host/self%> clause"); + "%<device%> or %<host%> or %<self%> clause"); return; } diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index ba1a8d8..33f5ead 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -13011,10 +13011,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; |