From: Cesar Philippidis Date: Fri, 6 Nov 2015 02:03:48 +0000 (-0800) Subject: gimplify.c (gimplify_scan_omp_clauses): Add support for OMP_CLAUSE_TILE. X-Git-Url: https://git.libre-soc.org/?a=commitdiff_plain;h=7a5e4956cc026cba54159d5c764486ac4151db85;p=gcc.git gimplify.c (gimplify_scan_omp_clauses): Add support for OMP_CLAUSE_TILE. gcc/ * gimplify.c (gimplify_scan_omp_clauses): Add support for OMP_CLAUSE_TILE. Update handling of OMP_CLAUSE_INDEPENDENT. (gimplify_adjust_omp_clauses): Likewise. * omp-low.c (scan_sharing_clauses): Add support for OMP_CLAUSE_TILE. * tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_TILE. * tree-pretty-print.c (dump_omp_clause): Handle OMP_CLAUSE_TILE. * tree.c (omp_clause_num_ops): Add an entry for OMP_CLAUSE_TILE. (omp_clause_code_name): Likewise. (walk_tree_1): Handle OMP_CLAUSE_TILE. * tree.h (OMP_TILE_LIST): New macro. gcc/c-family/ * c-omp.c (c_oacc_split_loop_clauses): Make TILE, GANG, WORKER, VECTOR, AUTO, SEQ, INDEPENDENT and PRIVATE loop clauses. Associate REDUCTION clauses with parallel and kernels and loops. * c-pragma.h (enum pragma_omp_clause): Add entries for PRAGMA_OACC_CLAUSE_{INDEPENDENT,TILE,DEFAULT}. * pt.c (tsubst_omp_clauses): Add support for OMP_CLAUSE_{NUM_GANGS, NUM_WORKERS,VECTOR_LENGTH,GANG,WORKER,VECTOR,ASYNC,WAIT,TILE,AUTO, INDEPENDENT,SEQ}. (tsubst_expr): Add support for OMP_CLAUSE_{KERNELS,PARALLEL,LOOP}. gcc/c/ * 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. gcc/cp/ * parser.c (cp_parser_omp_clause_name): Add support for PRAGMA_OACC_CLAUSE_INDEPENDENT and PRAGMA_OACC_CLAUSE_TILE. (cp_parser_oacc_shape_clause): Allow pointer variables as gang static arguments. (cp_parser_oacc_clause_tile): New function. (cp_parser_omp_clause_default): Add is_oacc argument. Handle default(none) in OpenACC. (cp_parser_oacc_all_clauses): Add support for (cp_parser_omp_all_clauses): Update call to cp_parser_omp_clause_default. PRAGMA_OACC_CLAUSE_{DEFAULT,INDEPENDENT,TILE,PRIVATE,FIRSTPRIVATE}. (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}. (cp_parser_oacc_update): Update the error message for missing clauses. * semantics.c (finish_omp_clauses): Add support for OMP_CLAUSE_INDEPENDENT and OMP_CLAUSE_TILE. gcc/fortran/ * openmp.c (gfc_match_omp_clauses): Update support for the tile and default clauses in OpenACC. (gfc_match_oacc_update): Error when data clauses are supplied. (oacc_compatible_clauses): Delete. (resolve_omp_clauses): Give special care for OpenACC reductions. Also update error reporting for the tile clause. (resolve_oacc_loop_blocks): Update error reporting for the tile clause. * trans-openmp.c (gfc_trans_omp_clauses): Update OMP_CLAUSE_SEQ. Add OMP_CLAUSE_{AUTO,TILE} and add support the the gang static argument. (gfc_trans_oacc_combined_directive): Update the list of clauses which are split to acc loops. gcc/testsuite/ * c-c++-common/goacc/combined-directives.c: New test. * c-c++-common/goacc/loop-clauses.c: New test. * c-c++-common/goacc/tile.c: New test. * c-c++-common/goacc/loop-shape.c: Add test for pointer variable as gang static arguments. * c-c++-common/goacc/update-1.c: Adjust expected error message. * g++.dg/goacc/template.C: New test. * gfortran.dg/goacc/combined-directives.f90: New test. * gfortran.dg/goacc/default.f95: New test. * gfortran.dg/goacc/default_none.f95: New test. * gfortran.dg/goacc/firstprivate-1.f95: New test. * gfortran.dg/goacc/gang-static.f95: New test. * gfortran.dg/goacc/kernels-loop-inner.f95: New test. * gfortran.dg/goacc/kernels-loops-adjacent.f95: New test. * gfortran.dg/goacc/list.f95: Update test. * gfortran.dg/goacc/loop-2.f95: Likewise. * gfortran.dg/goacc/loop-4.f95: New test. * gfortran.dg/goacc/loop-5.f95: New test. * gfortran.dg/goacc/loop-6.f95: New test. * gfortran.dg/goacc/loop-tree-1.f90: Update test. * gfortran.dg/goacc/multi-clause.f90: New test. * gfortran.dg/goacc/parallel-tree.f95: Update test. * gfortran.dg/goacc/update.f95: New test. Co-Authored-By: James Norris Co-Authored-By: Nathan Sidwell Co-Authored-By: Thomas Schwinge Co-Authored-By: Tom de Vries From-SVN: r229832 --- diff --git a/gcc/ChangeLog b/gcc/ChangeLog index 94e597e025a..f01e8ec1263 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,19 @@ +2015-11-05 Cesar Philippidis + Thomas Schwinge + James Norris + + + * gimplify.c (gimplify_scan_omp_clauses): Add support for + OMP_CLAUSE_TILE. Update handling of OMP_CLAUSE_INDEPENDENT. + (gimplify_adjust_omp_clauses): Likewise. + * omp-low.c (scan_sharing_clauses): Add support for OMP_CLAUSE_TILE. + * tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_TILE. + * tree-pretty-print.c (dump_omp_clause): Handle OMP_CLAUSE_TILE. + * tree.c (omp_clause_num_ops): Add an entry for OMP_CLAUSE_TILE. + (omp_clause_code_name): Likewise. + (walk_tree_1): Handle OMP_CLAUSE_TILE. + * tree.h (OMP_TILE_LIST): New macro. + 2015-11-05 Martin Sebor PR c++/67942 diff --git a/gcc/c-family/ChangeLog b/gcc/c-family/ChangeLog index c21f50410fa..cc9a6420407 100644 --- a/gcc/c-family/ChangeLog +++ b/gcc/c-family/ChangeLog @@ -1,3 +1,17 @@ +2015-11-05 Cesar Philippidis + Thomas Schwinge + James Norris + + * c-omp.c (c_oacc_split_loop_clauses): Make TILE, GANG, WORKER, VECTOR, + AUTO, SEQ, INDEPENDENT and PRIVATE loop clauses. Associate REDUCTION + clauses with parallel and kernels and loops. + * c-pragma.h (enum pragma_omp_clause): Add entries for + PRAGMA_OACC_CLAUSE_{INDEPENDENT,TILE,DEFAULT}. + * pt.c (tsubst_omp_clauses): Add support for OMP_CLAUSE_{NUM_GANGS, + NUM_WORKERS,VECTOR_LENGTH,GANG,WORKER,VECTOR,ASYNC,WAIT,TILE,AUTO, + INDEPENDENT,SEQ}. + (tsubst_expr): Add support for OMP_CLAUSE_{KERNELS,PARALLEL,LOOP}. + 2015-11-05 Martin Sebor PR c++/67942 diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c index a51611c8864..3e93b590fdb 100644 --- a/gcc/c-family/c-omp.c +++ b/gcc/c-family/c-omp.c @@ -867,7 +867,7 @@ c_omp_check_loop_iv_exprs (location_t stmt_loc, tree declv, tree decl, tree c_oacc_split_loop_clauses (tree clauses, tree *not_loop_clauses) { - tree next, loop_clauses; + tree next, loop_clauses, t; loop_clauses = *not_loop_clauses = NULL_TREE; for (; clauses ; clauses = next) @@ -876,12 +876,27 @@ c_oacc_split_loop_clauses (tree clauses, tree *not_loop_clauses) switch (OMP_CLAUSE_CODE (clauses)) { + /* Loop clauses. */ case OMP_CLAUSE_COLLAPSE: - case OMP_CLAUSE_REDUCTION: + case OMP_CLAUSE_TILE: + case OMP_CLAUSE_GANG: + case OMP_CLAUSE_WORKER: + case OMP_CLAUSE_VECTOR: + case OMP_CLAUSE_AUTO: + case OMP_CLAUSE_SEQ: + case OMP_CLAUSE_INDEPENDENT: + case OMP_CLAUSE_PRIVATE: OMP_CLAUSE_CHAIN (clauses) = loop_clauses; loop_clauses = clauses; break; + /* Reductions belong in both constructs. */ + case OMP_CLAUSE_REDUCTION: + t = copy_node (clauses); + OMP_CLAUSE_CHAIN (t) = loop_clauses; + loop_clauses = t; + + /* Parallel/kernels clauses. */ default: OMP_CLAUSE_CHAIN (clauses) = *not_loop_clauses; *not_loop_clauses = clauses; diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h index 69e7392eee3..953c4e31d04 100644 --- a/gcc/c-family/c-pragma.h +++ b/gcc/c-family/c-pragma.h @@ -153,6 +153,7 @@ enum pragma_omp_clause { PRAGMA_OACC_CLAUSE_DEVICEPTR, PRAGMA_OACC_CLAUSE_GANG, PRAGMA_OACC_CLAUSE_HOST, + PRAGMA_OACC_CLAUSE_INDEPENDENT, PRAGMA_OACC_CLAUSE_NUM_GANGS, PRAGMA_OACC_CLAUSE_NUM_WORKERS, PRAGMA_OACC_CLAUSE_PRESENT, @@ -162,6 +163,7 @@ enum pragma_omp_clause { PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE, PRAGMA_OACC_CLAUSE_SELF, PRAGMA_OACC_CLAUSE_SEQ, + PRAGMA_OACC_CLAUSE_TILE, PRAGMA_OACC_CLAUSE_VECTOR, PRAGMA_OACC_CLAUSE_VECTOR_LENGTH, PRAGMA_OACC_CLAUSE_WAIT, @@ -169,6 +171,7 @@ enum pragma_omp_clause { PRAGMA_OACC_CLAUSE_COLLAPSE = PRAGMA_OMP_CLAUSE_COLLAPSE, PRAGMA_OACC_CLAUSE_COPYIN = PRAGMA_OMP_CLAUSE_COPYIN, PRAGMA_OACC_CLAUSE_DEVICE = PRAGMA_OMP_CLAUSE_DEVICE, + PRAGMA_OACC_CLAUSE_DEFAULT = PRAGMA_OMP_CLAUSE_DEFAULT, PRAGMA_OACC_CLAUSE_FIRSTPRIVATE = PRAGMA_OMP_CLAUSE_FIRSTPRIVATE, PRAGMA_OACC_CLAUSE_IF = PRAGMA_OMP_CLAUSE_IF, PRAGMA_OACC_CLAUSE_PRIVATE = PRAGMA_OMP_CLAUSE_PRIVATE, diff --git a/gcc/c/ChangeLog b/gcc/c/ChangeLog index 7839024cc88..c734a81f45c 100644 --- a/gcc/c/ChangeLog +++ b/gcc/c/ChangeLog @@ -1,3 +1,26 @@ +2015-11-05 Cesar Philippidis + Thomas Schwinge + James Norris + + * 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 PR c/68090 diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index e2f55b354d1..23d0107eaf7 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 % or %"); + if (is_oacc) + c_parser_error (parser, "expected %"); + else + c_parser_error (parser, "expected % or %"); } 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, "% 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,"% 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 " - "% or % clause"); + "% or % or % clause"); return; } diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index ba1a8d862a6..33f5ead829f 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; diff --git a/gcc/cp/ChangeLog b/gcc/cp/ChangeLog index 1741fa2cfc0..e780c868b62 100644 --- a/gcc/cp/ChangeLog +++ b/gcc/cp/ChangeLog @@ -1,3 +1,27 @@ +2015-11-05 Cesar Philippidis + Thomas Schwinge + James Norris + + * parser.c (cp_parser_omp_clause_name): Add support for + PRAGMA_OACC_CLAUSE_INDEPENDENT and PRAGMA_OACC_CLAUSE_TILE. + (cp_parser_oacc_shape_clause): Allow pointer variables as gang static + arguments. + (cp_parser_oacc_clause_tile): New function. + (cp_parser_omp_clause_default): Add is_oacc argument. Handle + default(none) in OpenACC. + (cp_parser_oacc_all_clauses): Add support for + (cp_parser_omp_all_clauses): Update call to + cp_parser_omp_clause_default. + PRAGMA_OACC_CLAUSE_{DEFAULT,INDEPENDENT,TILE,PRIVATE,FIRSTPRIVATE}. + (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}. + (cp_parser_oacc_update): Update the error message for missing clauses. + * semantics.c (finish_omp_clauses): Add support for + OMP_CLAUSE_INDEPENDENT and OMP_CLAUSE_TILE. + 2015-11-05 Martin Sebor PR c++/67942 diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index d439c064fe3..c6f57297e72 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -29125,6 +29125,8 @@ cp_parser_omp_clause_name (cp_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; @@ -29219,6 +29221,8 @@ cp_parser_omp_clause_name (cp_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; @@ -29684,7 +29688,10 @@ cp_parser_oacc_shape_clause (cp_parser *parser, omp_clause_code kind, } /* Check for the '*' argument. */ - if (cp_lexer_next_token_is (lexer, CPP_MULT)) + if (cp_lexer_next_token_is (lexer, CPP_MULT) + && (cp_lexer_nth_token_is (parser->lexer, 2, CPP_COMMA) + || cp_lexer_nth_token_is (parser->lexer, 2, + CPP_CLOSE_PAREN))) { cp_lexer_consume_token (lexer); ops[idx] = integer_minus_one_node; @@ -29752,6 +29759,52 @@ cp_parser_oacc_shape_clause (cp_parser *parser, omp_clause_code kind, return list; } +/* OpenACC 2.0: + tile ( size-expr-list ) */ + +static tree +cp_parser_oacc_clause_tile (cp_parser *parser, location_t clause_loc, tree list) +{ + tree c, expr = error_mark_node; + tree tile = NULL_TREE; + + check_no_duplicate_clause (list, OMP_CLAUSE_TILE, "tile", clause_loc); + + if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN)) + return list; + + do + { + if (cp_lexer_next_token_is (parser->lexer, CPP_MULT) + && (cp_lexer_nth_token_is (parser->lexer, 2, CPP_COMMA) + || cp_lexer_nth_token_is (parser->lexer, 2, CPP_CLOSE_PAREN))) + { + cp_lexer_consume_token (parser->lexer); + expr = integer_minus_one_node; + } + else + expr = cp_parser_assignment_expression (parser, NULL, false, false); + + if (expr == error_mark_node) + return list; + + tile = tree_cons (NULL_TREE, expr, tile); + + if (cp_lexer_next_token_is (parser->lexer, CPP_COMMA)) + cp_lexer_consume_token (parser->lexer); + } + while (cp_lexer_next_token_is_not (parser->lexer, CPP_CLOSE_PAREN)); + + /* Consume the trailing ')'. */ + cp_lexer_consume_token (parser->lexer); + + c = build_omp_clause (clause_loc, OMP_CLAUSE_TILE); + tile = nreverse (tile); + OMP_CLAUSE_TILE_LIST (c) = tile; + OMP_CLAUSE_CHAIN (c) = list; + return c; +} + /* OpenACC 2.0 Parse wait clause or directive parameters. */ @@ -29859,10 +29912,14 @@ cp_parser_omp_clause_collapse (cp_parser *parser, tree list, location_t location } /* OpenMP 2.5: - default ( shared | none ) */ + default ( shared | none ) + + OpenACC 2.0 + default (none) */ static tree -cp_parser_omp_clause_default (cp_parser *parser, tree list, location_t location) +cp_parser_omp_clause_default (cp_parser *parser, tree list, + location_t location, bool is_oacc) { enum omp_clause_default_kind kind = OMP_CLAUSE_DEFAULT_UNSPECIFIED; tree c; @@ -29883,7 +29940,7 @@ cp_parser_omp_clause_default (cp_parser *parser, tree list, location_t location) break; case 's': - if (strcmp ("shared", p) != 0) + if (strcmp ("shared", p) != 0 || is_oacc) goto invalid_kind; kind = OMP_CLAUSE_DEFAULT_SHARED; break; @@ -29897,7 +29954,10 @@ cp_parser_omp_clause_default (cp_parser *parser, tree list, location_t location) else { invalid_kind: - cp_parser_error (parser, "expected % or %"); + if (is_oacc) + cp_parser_error (parser, "expected %"); + else + cp_parser_error (parser, "expected % or %"); } if (!cp_parser_require (parser, CPP_CLOSE_PAREN, RT_CLOSE_PAREN)) @@ -31444,6 +31504,10 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses); c_name = "delete"; break; + case PRAGMA_OMP_CLAUSE_DEFAULT: + clauses = cp_parser_omp_clause_default (parser, clauses, here, true); + c_name = "default"; + break; case PRAGMA_OACC_CLAUSE_DEVICE: clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses); c_name = "device"; @@ -31452,6 +31516,11 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, clauses = cp_parser_oacc_data_clause_deviceptr (parser, clauses); c_name = "deviceptr"; break; + case PRAGMA_OACC_CLAUSE_FIRSTPRIVATE: + clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_FIRSTPRIVATE, + clauses); + c_name = "firstprivate"; + break; case PRAGMA_OACC_CLAUSE_GANG: c_name = "gang"; clauses = cp_parser_oacc_shape_clause (parser, OMP_CLAUSE_GANG, @@ -31465,6 +31534,12 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, clauses = cp_parser_omp_clause_if (parser, clauses, here, false); c_name = "if"; break; + case PRAGMA_OACC_CLAUSE_INDEPENDENT: + clauses = cp_parser_oacc_simple_clause (parser, + OMP_CLAUSE_INDEPENDENT, + clauses, here); + c_name = "independent"; + break; case PRAGMA_OACC_CLAUSE_NUM_GANGS: code = OMP_CLAUSE_NUM_GANGS; c_name = "num_gangs"; @@ -31497,6 +31572,11 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses); c_name = "present_or_create"; break; + case PRAGMA_OACC_CLAUSE_PRIVATE: + clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_PRIVATE, + clauses); + c_name = "private"; + break; case PRAGMA_OACC_CLAUSE_REDUCTION: clauses = cp_parser_omp_clause_reduction (parser, clauses); c_name = "reduction"; @@ -31510,6 +31590,10 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, clauses, here); c_name = "seq"; break; + case PRAGMA_OACC_CLAUSE_TILE: + clauses = cp_parser_oacc_clause_tile (parser, here, clauses); + c_name = "tile"; + break; case PRAGMA_OACC_CLAUSE_VECTOR: c_name = "vector"; clauses = cp_parser_oacc_shape_clause (parser, OMP_CLAUSE_VECTOR, @@ -31598,7 +31682,7 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask, break; case PRAGMA_OMP_CLAUSE_DEFAULT: clauses = cp_parser_omp_clause_default (parser, clauses, - token->location); + token->location, false); c_name = "default"; break; case PRAGMA_OMP_CLAUSE_FINAL: @@ -34493,12 +34577,15 @@ cp_parser_oacc_enter_exit_data (cp_parser *parser, cp_token *pragma_tok, #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_VECTOR) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WORKER) \ | (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 cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name, @@ -34543,6 +34630,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, 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) \ @@ -34558,7 +34646,9 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, 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_FIRSTPRIVATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS) \ @@ -34567,6 +34657,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name, | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRIVATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) @@ -34642,7 +34733,7 @@ cp_parser_oacc_update (cp_parser *parser, cp_token *pragma_tok) { error_at (pragma_tok->location, "%<#pragma acc update%> must contain at least one " - "% or % clause"); + "% or % or % clause"); return NULL_TREE; } diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c index 9e3bd2d261c..45eda3a6d72 100644 --- a/gcc/cp/pt.c +++ b/gcc/cp/pt.c @@ -14395,6 +14395,13 @@ tsubst_omp_clauses (tree clauses, bool declare_simd, bool allow_fields, case OMP_CLAUSE_PRIORITY: case OMP_CLAUSE_ORDERED: case OMP_CLAUSE_HINT: + case OMP_CLAUSE_NUM_GANGS: + case OMP_CLAUSE_NUM_WORKERS: + case OMP_CLAUSE_VECTOR_LENGTH: + case OMP_CLAUSE_WORKER: + case OMP_CLAUSE_VECTOR: + case OMP_CLAUSE_ASYNC: + case OMP_CLAUSE_WAIT: OMP_CLAUSE_OPERAND (nc, 0) = tsubst_expr (OMP_CLAUSE_OPERAND (oc, 0), args, complain, in_decl, /*integral_constant_expression_p=*/false); @@ -14419,6 +14426,7 @@ tsubst_omp_clauses (tree clauses, bool declare_simd, bool allow_fields, = tsubst_omp_clause_decl (OMP_CLAUSE_DECL (oc), args, complain, in_decl); break; + case OMP_CLAUSE_GANG: case OMP_CLAUSE_ALIGNED: OMP_CLAUSE_DECL (nc) = tsubst_omp_clause_decl (OMP_CLAUSE_DECL (oc), args, complain, @@ -14461,6 +14469,22 @@ tsubst_omp_clauses (tree clauses, bool declare_simd, bool allow_fields, case OMP_CLAUSE_THREADS: case OMP_CLAUSE_SIMD: case OMP_CLAUSE_DEFAULTMAP: + case OMP_CLAUSE_INDEPENDENT: + case OMP_CLAUSE_AUTO: + case OMP_CLAUSE_SEQ: + break; + case OMP_CLAUSE_TILE: + { + tree lnc, loc; + for (lnc = OMP_CLAUSE_TILE_LIST (nc), + loc = OMP_CLAUSE_TILE_LIST (oc); + loc; + loc = TREE_CHAIN (loc), lnc = TREE_CHAIN (lnc)) + { + TREE_VALUE (lnc) = tsubst_expr (TREE_VALUE (loc), args, + complain, in_decl, false); + } + } break; default: gcc_unreachable (); @@ -15231,6 +15255,15 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl, } break; + case OACC_KERNELS: + case OACC_PARALLEL: + tmp = tsubst_omp_clauses (OMP_CLAUSES (t), false, false, args, complain, + in_decl); + stmt = begin_omp_parallel (); + RECUR (OMP_BODY (t)); + finish_omp_construct (TREE_CODE (t), stmt, tmp); + break; + case OMP_PARALLEL: r = push_omp_privatization_clauses (OMP_PARALLEL_COMBINED (t)); tmp = tsubst_omp_clauses (OMP_PARALLEL_CLAUSES (t), false, true, @@ -15261,6 +15294,7 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl, case CILK_FOR: case OMP_DISTRIBUTE: case OMP_TASKLOOP: + case OACC_LOOP: { tree clauses, body, pre_body; tree declv = NULL_TREE, initv = NULL_TREE, condv = NULL_TREE; @@ -15269,7 +15303,8 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl, int i; r = push_omp_privatization_clauses (OMP_FOR_INIT (t) == NULL_TREE); - clauses = tsubst_omp_clauses (OMP_FOR_CLAUSES (t), false, true, + clauses = tsubst_omp_clauses (OMP_FOR_CLAUSES (t), false, + TREE_CODE (t) != OACC_LOOP, args, complain, in_decl); if (OMP_FOR_INIT (t) != NULL_TREE) { @@ -15339,9 +15374,11 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl, pop_omp_privatization_clauses (r); break; + case OACC_DATA: case OMP_TARGET_DATA: case OMP_TARGET: - tmp = tsubst_omp_clauses (OMP_CLAUSES (t), false, true, + tmp = tsubst_omp_clauses (OMP_CLAUSES (t), false, + TREE_CODE (t) != OACC_DATA, args, complain, in_decl); keep_next_level (true); stmt = begin_omp_structured_block (); @@ -15395,6 +15432,16 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl, add_stmt (t); break; + case OACC_ENTER_DATA: + case OACC_EXIT_DATA: + case OACC_UPDATE: + tmp = tsubst_omp_clauses (OMP_STANDALONE_CLAUSES (t), false, false, + args, complain, in_decl); + t = copy_node (t); + OMP_STANDALONE_CLAUSES (t) = tmp; + add_stmt (t); + break; + case OMP_ORDERED: tmp = tsubst_omp_clauses (OMP_ORDERED_CLAUSES (t), false, true, args, complain, in_decl); diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c index fe18b273be5..7702a418b5b 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -6855,9 +6855,47 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd) case OMP_CLAUSE_DEFAULTMAP: case OMP_CLAUSE__CILK_FOR_COUNT_: case OMP_CLAUSE_AUTO: + case OMP_CLAUSE_INDEPENDENT: case OMP_CLAUSE_SEQ: break; + case OMP_CLAUSE_TILE: + for (tree list = OMP_CLAUSE_TILE_LIST (c); !remove && list; + list = TREE_CHAIN (list)) + { + t = TREE_VALUE (list); + + if (t == error_mark_node) + remove = true; + else if (!type_dependent_expression_p (t) + && !INTEGRAL_TYPE_P (TREE_TYPE (t))) + { + error ("% value must be integral"); + remove = true; + } + else + { + t = mark_rvalue_use (t); + if (!processing_template_decl) + { + t = maybe_constant_value (t); + if (TREE_CODE (t) == INTEGER_CST + && tree_int_cst_sgn (t) != 1 + && t != integer_minus_one_node) + { + warning_at (OMP_CLAUSE_LOCATION (c), 0, + "% value must be positive"); + t = integer_one_node; + } + } + t = fold_build_cleanup_point_expr (TREE_TYPE (t), t); + } + + /* Update list item. */ + TREE_VALUE (list) = t; + } + break; + case OMP_CLAUSE_ORDERED: ordered_seen = true; break; diff --git a/gcc/fortran/ChangeLog b/gcc/fortran/ChangeLog index e2c8f715fd5..bd3fa7647a5 100644 --- a/gcc/fortran/ChangeLog +++ b/gcc/fortran/ChangeLog @@ -1,3 +1,17 @@ +2015-11-05 Cesar Philippidis + + * openmp.c (gfc_match_omp_clauses): Update support for the tile + and default clauses in OpenACC. + (gfc_match_oacc_update): Error when data clauses are supplied. + (oacc_compatible_clauses): Delete. + (resolve_omp_clauses): Give special care for OpenACC reductions. + Also update error reporting for the tile clause. + (resolve_oacc_loop_blocks): Update error reporting for the tile clause. + * trans-openmp.c (gfc_trans_omp_clauses): Update OMP_CLAUSE_SEQ. Add + OMP_CLAUSE_{AUTO,TILE} and add support the the gang static argument. + (gfc_trans_oacc_combined_directive): Update the list of clauses which + are split to acc loops. + 2015-11-05 Jakub Jelinek * types.def (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR): Remove. diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c index 929a739972d..a7c7a1927e3 100644 --- a/gcc/fortran/openmp.c +++ b/gcc/fortran/openmp.c @@ -703,6 +703,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask, OMP_MAP_FORCE_FROM)) continue; if ((mask & OMP_CLAUSE_TILE) + && !c->tile_list && match_oacc_expr_list ("tile (", &c->tile_list, true) == MATCH_YES) continue; if ((mask & OMP_CLAUSE_SEQ) && !c->seq @@ -856,12 +857,14 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask, if ((mask & OMP_CLAUSE_DEFAULT) && c->default_sharing == OMP_DEFAULT_UNKNOWN) { - if (gfc_match ("default ( shared )") == MATCH_YES) + if (gfc_match ("default ( none )") == MATCH_YES) + c->default_sharing = OMP_DEFAULT_NONE; + else if (openacc) + /* c->default_sharing = OMP_DEFAULT_UNKNOWN */; + else if (gfc_match ("default ( shared )") == MATCH_YES) c->default_sharing = OMP_DEFAULT_SHARED; else if (gfc_match ("default ( private )") == MATCH_YES) c->default_sharing = OMP_DEFAULT_PRIVATE; - else if (gfc_match ("default ( none )") == MATCH_YES) - c->default_sharing = OMP_DEFAULT_NONE; else if (gfc_match ("default ( firstprivate )") == MATCH_YES) c->default_sharing = OMP_DEFAULT_FIRSTPRIVATE; if (c->default_sharing != OMP_DEFAULT_UNKNOWN) @@ -1304,10 +1307,19 @@ match gfc_match_oacc_update (void) { gfc_omp_clauses *c; + locus here = gfc_current_locus; + if (gfc_match_omp_clauses (&c, OACC_UPDATE_CLAUSES, false, false, true) != MATCH_YES) return MATCH_ERROR; + if (!c->lists[OMP_LIST_MAP]) + { + gfc_error ("% must contain at least one " + "% or % or % clause at %L", &here); + return MATCH_ERROR; + } + new_st.op = EXEC_OACC_UPDATE; new_st.ext.omp_clauses = c; return MATCH_YES; @@ -2846,30 +2858,6 @@ resolve_omp_udr_clause (gfc_omp_namelist *n, gfc_namespace *ns, return copy; } -/* Returns true if clause in list 'list' is compatible with any of - of the clauses in lists [0..list-1]. E.g., a reduction variable may - appear in both reduction and private clauses, so this function - will return true in this case. */ - -static bool -oacc_compatible_clauses (gfc_omp_clauses *clauses, int list, - gfc_symbol *sym, bool openacc) -{ - gfc_omp_namelist *n; - - if (!openacc) - return false; - - if (list != OMP_LIST_REDUCTION) - return false; - - for (n = clauses->lists[OMP_LIST_FIRST]; n; n = n->next) - if (n->sym == sym) - return true; - - return false; -} - /* OpenMP directive resolving routines. */ static void @@ -2975,11 +2963,11 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, && list != OMP_LIST_DEPEND && (list != OMP_LIST_MAP || openacc) && list != OMP_LIST_FROM - && list != OMP_LIST_TO) + && list != OMP_LIST_TO + && (list != OMP_LIST_REDUCTION || !openacc)) for (n = omp_clauses->lists[list]; n; n = n->next) { - if (n->sym->mark && !oacc_compatible_clauses (omp_clauses, list, - n->sym, openacc)) + if (n->sym->mark) gfc_error ("Symbol %qs present on multiple clauses at %L", n->sym->name, &n->where); else @@ -3028,6 +3016,22 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, n->sym->mark = 1; } + /* OpenACC reductions. */ + if (openacc) + { + for (n = omp_clauses->lists[OMP_LIST_REDUCTION]; n; n = n->next) + n->sym->mark = 0; + + for (n = omp_clauses->lists[OMP_LIST_REDUCTION]; n; n = n->next) + { + if (n->sym->mark) + gfc_error ("Symbol %qs present on multiple clauses at %L", + n->sym->name, &n->where); + else + n->sym->mark = 1; + } + } + for (n = omp_clauses->lists[OMP_LIST_TO]; n; n = n->next) n->sym->mark = 0; for (n = omp_clauses->lists[OMP_LIST_FROM]; n; n = n->next) @@ -4528,22 +4532,8 @@ resolve_oacc_loop_blocks (gfc_code *code) if (code->ext.omp_clauses->vector) gfc_error ("Clause AUTO conflicts with VECTOR at %L", &code->loc); } - if (!code->ext.omp_clauses->tile_list) - { - if (code->ext.omp_clauses->gang) - { - if (code->ext.omp_clauses->worker) - gfc_error ("Clause GANG conflicts with WORKER at %L", &code->loc); - if (code->ext.omp_clauses->vector) - gfc_error ("Clause GANG conflicts with VECTOR at %L", &code->loc); - } - if (code->ext.omp_clauses->worker) - if (code->ext.omp_clauses->vector) - gfc_error ("Clause WORKER conflicts with VECTOR at %L", &code->loc); - } - else if (code->ext.omp_clauses->gang - && code->ext.omp_clauses->worker - && code->ext.omp_clauses->vector) + if (code->ext.omp_clauses->tile_list && code->ext.omp_clauses->gang + && code->ext.omp_clauses->worker && code->ext.omp_clauses->vector) gfc_error ("Tiled loop cannot be parallelized across gangs, workers and " "vectors at the same time at %L", &code->loc); @@ -4564,10 +4554,21 @@ resolve_oacc_loop_blocks (gfc_code *code) { num++; if (el->expr == NULL) - continue; - resolve_oacc_positive_int_expr (el->expr, "TILE"); - if (el->expr->expr_type != EXPR_CONSTANT) - gfc_error ("TILE requires constant expression at %L", &code->loc); + { + /* NULL expressions are used to represent '*' arguments. + Convert those to a -1 expressions. */ + el->expr = gfc_get_constant_expr (BT_INTEGER, + gfc_default_integer_kind, + &code->loc); + mpz_set_si (el->expr->value.integer, -1); + } + else + { + resolve_oacc_positive_int_expr (el->expr, "TILE"); + if (el->expr->expr_type != EXPR_CONSTANT) + gfc_error ("TILE requires constant expression at %L", + &code->loc); + } } resolve_oacc_nested_loops (code, code->block->next, num, "tiled"); } diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index 5f4c382a5dd..9f0d5332a8e 100644 --- a/gcc/fortran/trans-openmp.c +++ b/gcc/fortran/trans-openmp.c @@ -2534,8 +2534,12 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, } if (clauses->seq) { - c = build_omp_clause (where.lb->location, OMP_CLAUSE_ORDERED); - OMP_CLAUSE_ORDERED_EXPR (c) = NULL_TREE; + c = build_omp_clause (where.lb->location, OMP_CLAUSE_SEQ); + omp_clauses = gfc_trans_add_clause (c, omp_clauses); + } + if (clauses->par_auto) + { + c = build_omp_clause (where.lb->location, OMP_CLAUSE_AUTO); omp_clauses = gfc_trans_add_clause (c, omp_clauses); } if (clauses->independent) @@ -2579,6 +2583,21 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, OMP_CLAUSE_VECTOR_LENGTH_EXPR (c) = vector_length_var; omp_clauses = gfc_trans_add_clause (c, omp_clauses); } + if (clauses->tile_list) + { + vec *tvec; + gfc_expr_list *el; + + vec_alloc (tvec, 4); + + for (el = clauses->tile_list; el; el = el->next) + vec_safe_push (tvec, gfc_convert_expr_to_tree (block, el->expr)); + + c = build_omp_clause (where.lb->location, OMP_CLAUSE_TILE); + OMP_CLAUSE_TILE_LIST (c) = build_tree_list_vec (tvec); + omp_clauses = gfc_trans_add_clause (c, omp_clauses); + tvec->truncate (0); + } if (clauses->vector) { if (clauses->vector_expr) @@ -2618,7 +2637,17 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, tree gang_var = gfc_convert_expr_to_tree (block, clauses->gang_expr); c = build_omp_clause (where.lb->location, OMP_CLAUSE_GANG); - OMP_CLAUSE_GANG_EXPR (c) = gang_var; + if (clauses->gang_static) + OMP_CLAUSE_GANG_STATIC_EXPR (c) = gang_var; + else + OMP_CLAUSE_GANG_EXPR (c) = gang_var; + omp_clauses = gfc_trans_add_clause (c, omp_clauses); + } + else if (clauses->gang_static) + { + /* This corresponds to gang (static: *). */ + c = build_omp_clause (where.lb->location, OMP_CLAUSE_GANG); + OMP_CLAUSE_GANG_STATIC_EXPR (c) = integer_minus_one_node; omp_clauses = gfc_trans_add_clause (c, omp_clauses); } else @@ -3449,16 +3478,33 @@ gfc_trans_oacc_combined_directive (gfc_code *code) sizeof (construct_clauses)); loop_clauses.collapse = construct_clauses.collapse; loop_clauses.gang = construct_clauses.gang; + loop_clauses.gang_expr = construct_clauses.gang_expr; + loop_clauses.gang_static = construct_clauses.gang_static; loop_clauses.vector = construct_clauses.vector; + loop_clauses.vector_expr = construct_clauses.vector_expr; loop_clauses.worker = construct_clauses.worker; + loop_clauses.worker_expr = construct_clauses.worker_expr; loop_clauses.seq = construct_clauses.seq; + loop_clauses.par_auto = construct_clauses.par_auto; loop_clauses.independent = construct_clauses.independent; - construct_clauses.collapse = 0; + loop_clauses.tile_list = construct_clauses.tile_list; + loop_clauses.lists[OMP_LIST_PRIVATE] + = construct_clauses.lists[OMP_LIST_PRIVATE]; + loop_clauses.lists[OMP_LIST_REDUCTION] + = construct_clauses.lists[OMP_LIST_REDUCTION]; construct_clauses.gang = false; + construct_clauses.gang_expr = NULL; + construct_clauses.gang_static = false; construct_clauses.vector = false; + construct_clauses.vector_expr = NULL; construct_clauses.worker = false; + construct_clauses.worker_expr = NULL; construct_clauses.seq = false; + construct_clauses.par_auto = false; + construct_clauses.independent = false; construct_clauses.independent = false; + construct_clauses.tile_list = NULL; + construct_clauses.lists[OMP_LIST_PRIVATE] = NULL; oacc_clauses = gfc_trans_omp_clauses (&block, &construct_clauses, code->loc); } diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 319e200f34b..fa348585a24 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -7141,9 +7141,18 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, remove = true; break; + case OMP_CLAUSE_TILE: + for (tree list = OMP_CLAUSE_TILE_LIST (c); !remove && list; + list = TREE_CHAIN (list)) + { + if (gimplify_expr (&TREE_VALUE (list), pre_p, NULL, + is_gimple_val, fb_rvalue) == GS_ERROR) + remove = true; + } + break; + case OMP_CLAUSE_DEVICE_RESIDENT: case OMP_CLAUSE_USE_DEVICE: - case OMP_CLAUSE_INDEPENDENT: remove = true; break; @@ -7153,6 +7162,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, case OMP_CLAUSE_COLLAPSE: case OMP_CLAUSE_AUTO: case OMP_CLAUSE_SEQ: + case OMP_CLAUSE_INDEPENDENT: case OMP_CLAUSE_MERGEABLE: case OMP_CLAUSE_PROC_BIND: case OMP_CLAUSE_SAFELEN: @@ -7663,6 +7673,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, tree *list_p, case OMP_CLAUSE_VECTOR: case OMP_CLAUSE_AUTO: case OMP_CLAUSE_SEQ: + case OMP_CLAUSE_TILE: break; default: diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 86cc65500af..5ffb2765059 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -2129,6 +2129,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_GANG: case OMP_CLAUSE_WORKER: case OMP_CLAUSE_VECTOR: + case OMP_CLAUSE_TILE: break; case OMP_CLAUSE_ALIGNED: @@ -2309,6 +2310,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_GANG: case OMP_CLAUSE_WORKER: case OMP_CLAUSE_VECTOR: + case OMP_CLAUSE_TILE: break; case OMP_CLAUSE_DEVICE_RESIDENT: diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog index eb518088d91..c074e7b55e0 100644 --- a/gcc/testsuite/ChangeLog +++ b/gcc/testsuite/ChangeLog @@ -1,3 +1,32 @@ +2015-11-05 Cesar Philippidis + Tom de Vries + Nathan Sidwell + Thomas Schwinge + + * c-c++-common/goacc/combined-directives.c: New test. + * c-c++-common/goacc/loop-clauses.c: New test. + * c-c++-common/goacc/tile.c: New test. + * c-c++-common/goacc/loop-shape.c: Add test for pointer variable + as gang static arguments. + * c-c++-common/goacc/update-1.c: Adjust expected error message. + * g++.dg/goacc/template.C: New test. + * gfortran.dg/goacc/combined-directives.f90: New test. + * gfortran.dg/goacc/default.f95: New test. + * gfortran.dg/goacc/default_none.f95: New test. + * gfortran.dg/goacc/firstprivate-1.f95: New test. + * gfortran.dg/goacc/gang-static.f95: New test. + * gfortran.dg/goacc/kernels-loop-inner.f95: New test. + * gfortran.dg/goacc/kernels-loops-adjacent.f95: New test. + * gfortran.dg/goacc/list.f95: Update test. + * gfortran.dg/goacc/loop-2.f95: Likewise. + * gfortran.dg/goacc/loop-4.f95: New test. + * gfortran.dg/goacc/loop-5.f95: New test. + * gfortran.dg/goacc/loop-6.f95: New test. + * gfortran.dg/goacc/loop-tree-1.f90: Update test. + * gfortran.dg/goacc/multi-clause.f90: New test. + * gfortran.dg/goacc/parallel-tree.f95: Update test. + * gfortran.dg/goacc/update.f95: New test. + 2015-11-05 Martin Sebor PR c++/67942 diff --git a/gcc/testsuite/c-c++-common/goacc/combined-directives.c b/gcc/testsuite/c-c++-common/goacc/combined-directives.c new file mode 100644 index 00000000000..c3872851298 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/combined-directives.c @@ -0,0 +1,119 @@ +// { dg-do compile } +// { dg-options "-fopenacc -fdump-tree-gimple" } + +// This error is temporary. Remove when support is added for these clauses +// in the middle end. Also remove the comments from the reduction test +// after the FE learns that reduction variables may appear in data clauses too. +// { dg-prune-output "sorry, unimplemented" } + +void +test () +{ + int a[100], i, j, z; + + // acc parallel + + #pragma acc parallel loop collapse (2) + for (i = 0; i < 100; i++) + for (j = 0; j < 10; j++) + ; + + #pragma acc parallel loop gang + for (i = 0; i < 100; i++) + ; + + #pragma acc parallel loop worker + for (i = 0; i < 100; i++) + for (j = 0; j < 10; j++) + ; + + #pragma acc parallel loop vector + for (i = 0; i < 100; i++) + for (j = 0; j < 10; j++) + ; + + #pragma acc parallel loop seq + for (i = 0; i < 100; i++) + for (j = 0; j < 10; j++) + ; + + #pragma acc parallel loop auto + for (i = 0; i < 100; i++) + for (j = 0; j < 10; j++) + ; + + #pragma acc parallel loop tile (2, 3) + for (i = 0; i < 100; i++) + for (j = 0; j < 10; j++) + ; + + #pragma acc parallel loop independent + for (i = 0; i < 100; i++) + ; + + #pragma acc parallel loop private (z) + for (i = 0; i < 100; i++) + z = 0; + +// #pragma acc parallel loop reduction (+:z) copy (z) +// for (i = 0; i < 100; i++) +// ; + + // acc kernels + + #pragma acc kernels loop collapse (2) + for (i = 0; i < 100; i++) + for (j = 0; j < 10; j++) + ; + + #pragma acc kernels loop gang + for (i = 0; i < 100; i++) + ; + + #pragma acc kernels loop worker + for (i = 0; i < 100; i++) + for (j = 0; j < 10; j++) + ; + + #pragma acc kernels loop vector + for (i = 0; i < 100; i++) + for (j = 0; j < 10; j++) + ; + + #pragma acc kernels loop seq + for (i = 0; i < 100; i++) + for (j = 0; j < 10; j++) + ; + + #pragma acc kernels loop auto + for (i = 0; i < 100; i++) + for (j = 0; j < 10; j++) + ; + + #pragma acc kernels loop tile (2, 3) + for (i = 0; i < 100; i++) + for (j = 0; j < 10; j++) + ; + + #pragma acc kernels loop independent + for (i = 0; i < 100; i++) + ; + + #pragma acc kernels loop private (z) + for (i = 0; i < 100; i++) + z = 0; + +// #pragma acc kernels loop reduction (+:z) copy (z) +// for (i = 0; i < 100; i++) +// ; +} + +// { dg-final { scan-tree-dump-times "acc loop collapse.2. private.j. private.i" 2 "gimple" } } +// { dg-final { scan-tree-dump-times "acc loop gang" 2 "gimple" } } +// { dg-final { scan-tree-dump-times "acc loop worker" 2 "gimple" } } +// { dg-final { scan-tree-dump-times "acc loop vector" 2 "gimple" } } +// { dg-final { scan-tree-dump-times "acc loop seq" 2 "gimple" } } +// { dg-final { scan-tree-dump-times "acc loop auto" 2 "gimple" } } +// { dg-final { scan-tree-dump-times "acc loop tile.2, 3" 2 "gimple" } } +// { dg-final { scan-tree-dump-times "acc loop independent private.i" 2 "gimple" } } +// { dg-final { scan-tree-dump-times "private.z" 2 "gimple" } } diff --git a/gcc/testsuite/c-c++-common/goacc/loop-clauses.c b/gcc/testsuite/c-c++-common/goacc/loop-clauses.c new file mode 100644 index 00000000000..97b8786a290 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/loop-clauses.c @@ -0,0 +1,86 @@ +/* { dg-do compile } */ + +/* { dg-prune-output "sorry, unimplemented" } */ + +int +main () +{ + int i, j; + +#pragma acc parallel firstprivate (j) private (i) + { +#pragma acc loop seq + for (i = 0; i < 10; i++) + { } + } + +#pragma acc parallel default (none) + { +#pragma acc loop auto private (j) + for (i = 0; i < 10; i++) + { } +#pragma acc loop gang + for (i = 0; i < 10; i++) + { } +#pragma acc loop gang(static:5) + for (i = 0; i < 10; i++) + { } +#pragma acc loop gang(static:*) + for (i = 0; i < 10; i++) + { } +#pragma acc loop vector + for (i = 0; i < 10; i++) + { } +#pragma acc loop worker + for (i = 0; i < 10; i++) + { } +#pragma acc loop auto + for (i = 0; i < 10; i++) + { } +#pragma acc loop independent + for (i = 0; i < 10; i++) + { } +#pragma acc loop seq + for (i = 0; i < 10; i++) + { } +#pragma acc loop gang worker vector + for (i = 0; i < 10; i++) + { } + } + +#pragma acc kernels default (none) + { +#pragma acc loop auto + for (i = 0; i < 10; i++) + { } +#pragma acc loop gang (num:5) + for (i = 0; i < 10; i++) + { } +#pragma acc loop gang(static:5) + for (i = 0; i < 10; i++) + { } +#pragma acc loop gang(static:*) + for (i = 0; i < 10; i++) + { } +#pragma acc loop vector(length:10) + for (i = 0; i < 10; i++) + { } +#pragma acc loop worker(num:5) + for (i = 0; i < 10; i++) + { } +#pragma acc loop auto + for (i = 0; i < 10; i++) + { } +#pragma acc loop independent + for (i = 0; i < 10; i++) + { } +#pragma acc loop seq + for (i = 0; i < 10; i++) + { } +#pragma acc loop gang worker vector + for (i = 0; i < 10; i++) + { } + } + + return 0; +} diff --git a/gcc/testsuite/c-c++-common/goacc/loop-shape.c b/gcc/testsuite/c-c++-common/goacc/loop-shape.c index b6d3156074d..9708f7bf5eb 100644 --- a/gcc/testsuite/c-c++-common/goacc/loop-shape.c +++ b/gcc/testsuite/c-c++-common/goacc/loop-shape.c @@ -8,6 +8,7 @@ int main () int i; int v = 32, w = 19; int length = 1, num = 5; + int *abc; /* Valid uses. */ @@ -199,12 +200,12 @@ int main () ; #pragma acc kernels - #pragma acc loop gang(static: * abc) /* { dg-error "expected '.' before" } */ + #pragma acc loop gang(static: * abc) for (i = 0; i < 10; i++) ; #pragma acc kernels - #pragma acc loop gang(static:*num:1) /* { dg-error "expected '.' before" } */ + #pragma acc loop gang(static:*num:1) /* { dg-error "" } */ for (i = 0; i < 10; i++) ; diff --git a/gcc/testsuite/c-c++-common/goacc/tile.c b/gcc/testsuite/c-c++-common/goacc/tile.c new file mode 100644 index 00000000000..2a81427f053 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/tile.c @@ -0,0 +1,73 @@ +/* { dg-do compile } */ + +int +main () +{ + int i, *a, b; + +#pragma acc parallel loop tile (10) + for (i = 0; i < 100; i++) + ; + +#pragma acc parallel loop tile (*) + for (i = 0; i < 100; i++) + ; + +#pragma acc parallel loop tile (10, *) + for (i = 0; i < 100; i++) + ; + +#pragma acc parallel loop tile (10, *, i) + for (i = 0; i < 100; i++) + ; + +#pragma acc parallel loop tile // { dg-error "expected '\\\('" } + for (i = 0; i < 100; i++) + ; + +#pragma acc parallel loop tile () // { dg-error "" } + for (i = 0; i < 100; i++) + ; + +#pragma acc parallel loop tile (,1) // { dg-error "" } + for (i = 0; i < 100; i++) + ; + +#pragma acc parallel loop tile (,,) // { dg-error "" } + for (i = 0; i < 100; i++) + ; + +#pragma acc parallel loop tile (1.1) // { dg-error "'tile' value must be integral" } + for (i = 0; i < 100; i++) + ; + +#pragma acc parallel loop tile (-3) // { dg-warning "'tile' value must be positive" } + for (i = 0; i < 100; i++) + ; + +#pragma acc parallel loop tile (10,-3) // { dg-warning "'tile' value must be positive" } + for (i = 0; i < 100; i++) + ; + +#pragma acc parallel loop tile (-100,10,5) // { dg-warning "'tile' value must be positive" } + for (i = 0; i < 100; i++) + ; + +#pragma acc parallel loop tile (1,2.0,true) // { dg-error "" } + for (i = 0; i < 100; i++) + ; + +#pragma acc parallel loop tile (*a, 1) + for (i = 0; i < 100; i++) + ; + +#pragma acc parallel loop tile (1, *a, b) + for (i = 0; i < 100; i++) + ; + +#pragma acc parallel loop tile (b, 1, *a) + for (i = 0; i < 100; i++) + ; + + return 0; +} diff --git a/gcc/testsuite/c-c++-common/goacc/update-1.c b/gcc/testsuite/c-c++-common/goacc/update-1.c index 97e93794934..701ef363643 100644 --- a/gcc/testsuite/c-c++-common/goacc/update-1.c +++ b/gcc/testsuite/c-c++-common/goacc/update-1.c @@ -1,7 +1,7 @@ void f (void) { -#pragma acc update /* { dg-error "'#pragma acc update' must contain at least one 'device' or 'host/self' clause" } */ +#pragma acc update /* { dg-error "'#pragma acc update' must contain at least one 'device' or 'host' or 'self' clause" } */ int i = 0; int a[10]; diff --git a/gcc/testsuite/g++.dg/goacc/template.C b/gcc/testsuite/g++.dg/goacc/template.C new file mode 100644 index 00000000000..f7a717bf7ed --- /dev/null +++ b/gcc/testsuite/g++.dg/goacc/template.C @@ -0,0 +1,141 @@ +// This error is temporary. Remove when support is added for these clauses +// in the middle end. Also remove the comments from the reduction test +// after the FE learns that reduction variables may appear in data clauses too. +// { dg-prune-output "sorry, unimplemented" } + +#pragma acc routine +template T +accDouble(int val) +{ + return val * 2; +} + +template T +oacc_parallel_copy (T a) +{ + T b = 0; + char w = 1; + int x = 2; + float y = 3; + double z = 4; + +#pragma acc parallel num_gangs (a) num_workers (a) vector_length (a) default (none) copyout (b) copyin (a) + { + b = a; + } + +#pragma acc parallel num_gangs (a) copy (w, x, y, z) + { + w = accDouble(w); + x = accDouble(x); + y = accDouble(y); + z = accDouble(z); + } + +#pragma acc parallel num_gangs (a) if (1) + { +#pragma acc loop auto tile (a, 3) + for (int i = 0; i < a; i++) + for (int j = 0; j < 5; j++) + b = a; + +#pragma acc loop seq + for (int i = 0; i < a; i++) + b = a; + } + + T c; + +#pragma acc parallel num_workers (10) + { +#pragma acc atomic capture + c = b++; + +#pragma atomic update + c++; + +#pragma acc atomic read + b = a; + +#pragma acc atomic write + b = a; + } + +//#pragma acc parallel reduction (+:c) +// { +// c = 1; +// } + +#pragma acc data if (1) copy (b) + { + #pragma acc parallel + { + b = a; + } + } + +#pragma acc enter data copyin (b) +#pragma acc parallel present (b) + { + b = a; + } + +#pragma acc update host (b) +#pragma acc update self (b) +#pragma acc update device (b) +#pragma acc exit data delete (b) + + return b; +} + +template T +oacc_kernels_copy (T a) +{ + T b = 0; + T c = 0; + char w = 1; + int x = 2; + float y = 3; + double z = 4; + +#pragma acc kernels copy (w, x, y, z) + { + w = accDouble(w); + x = accDouble(x); + y = accDouble(y); + z = accDouble(z); + } + +#pragma acc kernels copyout (b) copyin (a) + b = a; + +//#pragma acc kernels loop reduction (+:c) +// for (int i = 0; i < 10; i++) +// { +// c = 1; +// } + +#pragma acc data if (1) copy (b) + { + #pragma acc kernels + { + b = a; + } + } + +#pragma acc enter data copyin (b) +#pragma acc kernels present (b) + { + b = a; + } + return b; +} + +int +main () +{ + int b = oacc_parallel_copy (5); + int c = oacc_kernels_copy (5); + + return b + c; +} diff --git a/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90 b/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90 new file mode 100644 index 00000000000..69775257cbd --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90 @@ -0,0 +1,155 @@ +! Exercise combined OpenACC directives. + +! { dg-do compile } +! { dg-options "-fopenacc -fdump-tree-gimple" } + +! This error is temporary. Remove when support is added for these clauses +! in the middle end. +! { dg-prune-output "sorry, unimplemented" } + +! Update the reduction tests. + +subroutine test + implicit none + integer a(100), i, j, z + + ! PARALLEL + + !$acc parallel loop collapse (2) + do i = 1, 100 + do j = 1, 10 + end do + end do + !$acc end parallel loop + + !$acc parallel loop gang + do i = 1, 100 + end do + !$acc end parallel loop + + !$acc parallel loop worker + do i = 1, 100 + do j = 1, 10 + end do + end do + !$acc end parallel loop + + !$acc parallel loop vector + do i = 1, 100 + do j = 1, 10 + end do + end do + !$acc end parallel loop + + !$acc parallel loop seq + do i = 1, 100 + do j = 1, 10 + end do + end do + !$acc end parallel loop + + !$acc parallel loop auto + do i = 1, 100 + do j = 1, 10 + end do + end do + !$acc end parallel loop + + !$acc parallel loop tile (2, 3) + do i = 1, 100 + do j = 1, 10 + end do + end do + !$acc end parallel loop + + !$acc parallel loop independent + do i = 1, 100 + end do + !$acc end parallel loop + + !$acc parallel loop private (z) + do i = 1, 100 + z = 0 + end do + !$acc end parallel loop + +! !$acc parallel loop reduction (+:z) copy (z) +! do i = 1, 100 +! end do +! !$acc end parallel loop + + ! KERNELS + + !$acc kernels loop collapse (2) + do i = 1, 100 + do j = 1, 10 + end do + end do + !$acc end kernels loop + + !$acc kernels loop gang + do i = 1, 100 + end do + !$acc end kernels loop + + !$acc kernels loop worker + do i = 1, 100 + do j = 1, 10 + end do + end do + !$acc end kernels loop + + !$acc kernels loop vector + do i = 1, 100 + do j = 1, 10 + end do + end do + !$acc end kernels loop + + !$acc kernels loop seq + do i = 1, 100 + do j = 1, 10 + end do + end do + !$acc end kernels loop + + !$acc kernels loop auto + do i = 1, 100 + do j = 1, 10 + end do + end do + !$acc end kernels loop + + !$acc kernels loop tile (2, 3) + do i = 1, 100 + do j = 1, 10 + end do + end do + !$acc end kernels loop + + !$acc kernels loop independent + do i = 1, 100 + end do + !$acc end kernels loop + + !$acc kernels loop private (z) + do i = 1, 100 + z = 0 + end do + !$acc end kernels loop + +! !$acc kernels loop reduction (+:z) copy (z) +! do i = 1, 100 +! end do +! !$acc end kernels loop +end subroutine test + +! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. collapse.2." 2 "gimple" } } +! { dg-final { scan-tree-dump-times "acc loop private.i. gang" 2 "gimple" } } +! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. worker" 2 "gimple" } } +! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. vector" 2 "gimple" } } +! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. seq" 2 "gimple" } } +! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. auto" 2 "gimple" } } +! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. tile.2, 3" 2 "gimple" } } +! { dg-final { scan-tree-dump-times "acc loop private.i. independent" 2 "gimple" } } +! { dg-final { scan-tree-dump-times "private.z" 2 "gimple" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/default.f95 b/gcc/testsuite/gfortran.dg/goacc/default.f95 new file mode 100644 index 00000000000..c1fc52e6014 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/default.f95 @@ -0,0 +1,17 @@ +! { dg-do compile } + +program tile + integer i, j, a + + !$acc parallel default (shared) ! { dg-error "Unclassifiable OpenACC directive" } + !$acc end parallel ! { dg-error "Unexpected" } + + !$acc parallel default (private) ! { dg-error "Unclassifiable OpenACC directive" } + !$acc end parallel ! { dg-error "Unexpected" } + + !$acc parallel default (none) + !$acc end parallel + + !$acc parallel default (firstprivate) ! { dg-error "Unclassifiable OpenACC directive" } + !$acc end parallel ! { dg-error "Unexpected" } +end program tile diff --git a/gcc/testsuite/gfortran.dg/goacc/default_none.f95 b/gcc/testsuite/gfortran.dg/goacc/default_none.f95 new file mode 100644 index 00000000000..5ce66aeacf2 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/default_none.f95 @@ -0,0 +1,59 @@ +! Ensure that the internal array variables, offset, lbound, etc., don't +! trigger errors with default(none). + +! { dg-do compile } + +program main + implicit none + integer i + integer,parameter :: n = 100 + integer,allocatable :: a1(:), a2(:,:) + + allocate (a1 (n)) + allocate (a2 (-n:n,-n:n)) + a1 (:) = -1 + + !$acc parallel loop default(none) copy (a1(1:n)) + do i = 1,n + a1(i) = i + end do + !$acc end parallel loop + + call foo (a1) + call bar (a1, n) + call foobar (a2,n) + +contains + + subroutine foo (da1) + integer :: da1(n) + + !$acc parallel loop default(none) copy (da1(1:n)) + do i = 1,n + da1(i) = i*2 + end do + !$acc end parallel loop + end subroutine foo +end program main + +subroutine bar (da2,n) + integer :: n, da2(n) + integer i + + !$acc parallel loop default(none) copy (da2(1:n)) firstprivate(n) + do i = 1,n + da2(i) = i*3 + end do + !$acc end parallel loop +end subroutine bar + +subroutine foobar (da3,n) + integer :: n, da3(-n:n,-n:n) + integer i + + !$acc parallel loop default(none) copy (da3(-n:n,-n:n)) firstprivate(n) + do i = 1,n + da3(i,0) = i*3 + end do + !$acc end parallel loop +end subroutine foobar diff --git a/gcc/testsuite/gfortran.dg/goacc/firstprivate-1.f95 b/gcc/testsuite/gfortran.dg/goacc/firstprivate-1.f95 new file mode 100644 index 00000000000..fb92deedadf --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/firstprivate-1.f95 @@ -0,0 +1,11 @@ +! { dg-do compile } + +program test + integer a, b(100) + + !$acc parallel firstprivate (a, b) + !$acc end parallel + + !$acc parallel firstprivate (b(10:20)) ! { dg-error "Syntax error in OpenMP variable list" } + !$acc end parallel ! { dg-error "Unexpected !\\\$ACC END PARALLEL statement" } +end program test diff --git a/gcc/testsuite/gfortran.dg/goacc/gang-static.f95 b/gcc/testsuite/gfortran.dg/goacc/gang-static.f95 new file mode 100644 index 00000000000..4e46cf3cb41 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/gang-static.f95 @@ -0,0 +1,68 @@ +! { dg-do compile } +! { dg-additional-options "-fdump-tree-omplower" } + +program main + integer, parameter :: n = 100 + integer i, a(n), b(n) + + do i = 1, n + b(i) = i + end do + + !$acc parallel loop gang (static:*) num_gangs (10) + do i = 1, n + a(i) = b(i) + 0 + end do + !$acc end parallel loop + + call test (a, b, 0, n) + + !$acc parallel loop gang (static:1) num_gangs (10) + do i = 1, n + a(i) = b(i) + 1 + end do + !$acc end parallel loop + + call test (a, b, 1, n) + + !$acc parallel loop gang (static:2) num_gangs (10) + do i = 1, n + a(i) = b(i) + 2 + end do + !$acc end parallel loop + + call test (a, b, 2, n) + + !$acc parallel loop gang (static:5) num_gangs (10) + do i = 1, n + a(i) = b(i) + 5 + end do + !$acc end parallel loop + + call test (a, b, 5, n) + + !$acc parallel loop gang (static:20) num_gangs (10) + do i = 1, n + a(i) = b(i) + 20 + end do + !$acc end parallel loop + + call test (a, b, 20, n) + +end program main + +subroutine test (a, b, sarg, n) + integer n + integer a (n), b(n), sarg + integer i + + do i = 1, n + if (a(i) .ne. b(i) + sarg) call abort () + end do +end subroutine test + +! { dg-final { scan-tree-dump-times "gang\\(static:\\\*\\)" 1 "omplower" } } +! { dg-final { scan-tree-dump-times "gang\\(static:1\\)" 1 "omplower" } } +! { dg-final { scan-tree-dump-times "gang\\(static:2\\)" 1 "omplower" } } +! { dg-final { scan-tree-dump-times "gang\\(static:5\\)" 1 "omplower" } } +! { dg-final { scan-tree-dump-times "gang\\(static:20\\)" 1 "omplower" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-inner.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-inner.f95 new file mode 100644 index 00000000000..4db3a506c05 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-inner.f95 @@ -0,0 +1,23 @@ +! { dg-additional-options "-O2" } +! { dg-additional-options "-ftree-parallelize-loops=32" } + +program main + implicit none + + integer :: a(100,100), b(100,100) + integer :: i, j, d + + !$acc kernels + do i=1,100 + do j=1,100 + a(i,j) = 1 + b(i,j) = 2 + a(i,j) = a(i,j) + b(i,j) + end do + end do + !$acc end kernels + + d = sum(a) + + print *,d +end program main diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loops-adjacent.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loops-adjacent.f95 new file mode 100644 index 00000000000..fef3d106d4b --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loops-adjacent.f95 @@ -0,0 +1,19 @@ +! { dg-additional-options "-O2" } +! { dg-additional-options "-ftree-parallelize-loops=10" } + +program main + implicit none + + integer :: a(10000), b(10000) + integer :: d + + !$acc kernels + a = 1 + b = 2 + a = a + b + !$acc end kernels + + d = sum(a) + + print *,d +end program main diff --git a/gcc/testsuite/gfortran.dg/goacc/list.f95 b/gcc/testsuite/gfortran.dg/goacc/list.f95 index 94fdadd86db..a8006bcdd84 100644 --- a/gcc/testsuite/gfortran.dg/goacc/list.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/list.f95 @@ -5,7 +5,7 @@ program test implicit none integer :: i, j, k, l, a(10) - common /b/ j, k + common /b/ k real, pointer :: p1 => NULL() complex :: c, d(10) @@ -108,4 +108,4 @@ program test !$acc host_data use_device(p1) ! { dg-error "POINTER" } !$acc end host_data -end program test \ No newline at end of file +end program test diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-2.f95 b/gcc/testsuite/gfortran.dg/goacc/loop-2.f95 index f85691eb8e3..b5e6368a493 100644 --- a/gcc/testsuite/gfortran.dg/goacc/loop-2.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/loop-2.f95 @@ -66,7 +66,7 @@ program test !$acc loop seq worker ! { dg-error "conflicts with" } DO i = 1,10 ENDDO - !$acc loop gang worker ! { dg-error "conflicts with" } + !$acc loop gang worker DO i = 1,10 ENDDO @@ -94,10 +94,10 @@ program test !$acc loop seq vector ! { dg-error "conflicts with" } DO i = 1,10 ENDDO - !$acc loop gang vector ! { dg-error "conflicts with" } + !$acc loop gang vector DO i = 1,10 ENDDO - !$acc loop worker vector ! { dg-error "conflicts with" } + !$acc loop worker vector DO i = 1,10 ENDDO @@ -239,7 +239,7 @@ program test !$acc loop seq worker ! { dg-error "conflicts with" } DO i = 1,10 ENDDO - !$acc loop gang worker ! { dg-error "conflicts with" } + !$acc loop gang worker DO i = 1,10 ENDDO @@ -267,10 +267,10 @@ program test !$acc loop seq vector ! { dg-error "conflicts with" } DO i = 1,10 ENDDO - !$acc loop gang vector ! { dg-error "conflicts with" } + !$acc loop gang vector DO i = 1,10 ENDDO - !$acc loop worker vector ! { dg-error "conflicts with" } + !$acc loop worker vector DO i = 1,10 ENDDO @@ -392,7 +392,7 @@ program test !$acc kernels loop seq worker ! { dg-error "conflicts with" } DO i = 1,10 ENDDO - !$acc kernels loop gang worker ! { dg-error "conflicts with" } + !$acc kernels loop gang worker DO i = 1,10 ENDDO @@ -420,10 +420,10 @@ program test !$acc kernels loop seq vector ! { dg-error "conflicts with" } DO i = 1,10 ENDDO - !$acc kernels loop gang vector ! { dg-error "conflicts with" } + !$acc kernels loop gang vector DO i = 1,10 ENDDO - !$acc kernels loop worker vector ! { dg-error "conflicts with" } + !$acc kernels loop worker vector DO i = 1,10 ENDDO @@ -544,7 +544,7 @@ program test !$acc parallel loop seq worker ! { dg-error "conflicts with" } DO i = 1,10 ENDDO - !$acc parallel loop gang worker ! { dg-error "conflicts with" } + !$acc parallel loop gang worker DO i = 1,10 ENDDO @@ -572,10 +572,10 @@ program test !$acc parallel loop seq vector ! { dg-error "conflicts with" } DO i = 1,10 ENDDO - !$acc parallel loop gang vector ! { dg-error "conflicts with" } + !$acc parallel loop gang vector DO i = 1,10 ENDDO - !$acc parallel loop worker vector ! { dg-error "conflicts with" } + !$acc parallel loop worker vector DO i = 1,10 ENDDO @@ -646,4 +646,4 @@ program test !$acc parallel loop gang worker tile(*) DO i = 1,10 ENDDO -end \ No newline at end of file +end diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-4.f95 b/gcc/testsuite/gfortran.dg/goacc/loop-4.f95 new file mode 100644 index 00000000000..7c53c022de4 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/loop-4.f95 @@ -0,0 +1,7 @@ +! Ensure that loops not affiliated with acc compute regions cause an error. + +subroutine test1 + !$acc loop gang ! { dg-error "loop directive must be associated with an OpenACC compute region" } + DO i = 1,10 + ENDDO +end subroutine test1 diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-5.f95 b/gcc/testsuite/gfortran.dg/goacc/loop-5.f95 new file mode 100644 index 00000000000..5cbd975bea6 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/loop-5.f95 @@ -0,0 +1,363 @@ +! { dg-do compile } +! { dg-additional-options "-fmax-errors=100" } + +! { dg-prune-output "sorry, unimplemented" } +! { dg-prune-output "Error: work-sharing region" } + +program test + implicit none + integer :: i, j + + !$acc kernels + !$acc loop auto + DO i = 1,10 + ENDDO + !$acc loop gang + DO i = 1,10 + ENDDO + !$acc loop gang(5) + DO i = 1,10 + ENDDO + !$acc loop gang(num:5) + DO i = 1,10 + ENDDO + !$acc loop gang(static:5) + DO i = 1,10 + ENDDO + !$acc loop gang(static:*) + DO i = 1,10 + ENDDO + !$acc loop gang + DO i = 1,10 + !$acc loop vector + DO j = 1,10 + ENDDO + !$acc loop worker + DO j = 1,10 + ENDDO + ENDDO + + !$acc loop worker + DO i = 1,10 + ENDDO + !$acc loop worker(5) + DO i = 1,10 + ENDDO + !$acc loop worker(num:5) + DO i = 1,10 + ENDDO + !$acc loop worker + DO i = 1,10 + !$acc loop vector + DO j = 1,10 + ENDDO + ENDDO + !$acc loop gang worker + DO i = 1,10 + ENDDO + + !$acc loop vector + DO i = 1,10 + ENDDO + !$acc loop vector(5) + DO i = 1,10 + ENDDO + !$acc loop vector(length:5) + DO i = 1,10 + ENDDO + !$acc loop vector + DO i = 1,10 + ENDDO + !$acc loop gang vector + DO i = 1,10 + ENDDO + !$acc loop worker vector + DO i = 1,10 + ENDDO + + !$acc loop auto + DO i = 1,10 + ENDDO + + !$acc loop tile(1) + DO i = 1,10 + ENDDO + !$acc loop tile(2) + DO i = 1,10 + ENDDO + !$acc loop tile(6-2) + DO i = 1,10 + ENDDO + !$acc loop tile(6+2) + DO i = 1,10 + ENDDO + !$acc loop tile(*) + DO i = 1,10 + ENDDO + !$acc loop tile(*, 1) + DO i = 1,10 + DO j = 1,10 + ENDDO + ENDDO + !$acc loop tile(-1) ! { dg-warning "must be positive" } + do i = 1,10 + enddo + !$acc loop vector tile(*) + DO i = 1,10 + ENDDO + !$acc loop worker tile(*) + DO i = 1,10 + ENDDO + !$acc loop gang tile(*) + DO i = 1,10 + ENDDO + !$acc loop vector gang tile(*) + DO i = 1,10 + ENDDO + !$acc loop vector worker tile(*) + DO i = 1,10 + ENDDO + !$acc loop gang worker tile(*) + DO i = 1,10 + ENDDO + !$acc end kernels + + + !$acc parallel + !$acc loop tile(1) + DO i = 1,10 + ENDDO + !$acc loop tile(*) + DO i = 1,10 + ENDDO + !$acc loop tile(2) + DO i = 1,10 + DO j = 1,10 + ENDDO + ENDDO + !$acc loop tile(-1) ! { dg-warning "must be positive" } + do i = 1,10 + enddo + !$acc loop vector tile(*) + DO i = 1,10 + ENDDO + !$acc loop worker tile(*) + DO i = 1,10 + ENDDO + !$acc loop gang tile(*) + DO i = 1,10 + ENDDO + !$acc loop vector gang tile(*) + DO i = 1,10 + ENDDO + !$acc loop vector worker tile(*) + DO i = 1,10 + ENDDO + !$acc loop gang worker tile(*) + DO i = 1,10 + ENDDO + !$acc end parallel + + !$acc kernels loop auto + DO i = 1,10 + ENDDO + !$acc kernels loop gang + DO i = 1,10 + ENDDO + !$acc kernels loop gang(5) + DO i = 1,10 + ENDDO + !$acc kernels loop gang(num:5) + DO i = 1,10 + ENDDO + !$acc kernels loop gang(static:5) + DO i = 1,10 + ENDDO + !$acc kernels loop gang(static:*) + DO i = 1,10 + ENDDO + !$acc kernels loop gang + DO i = 1,10 + !$acc kernels loop gang ! { dg-error "OpenACC construct inside of non-OpenACC region" } + DO j = 1,10 + ENDDO + ENDDO + + !$acc kernels loop worker + DO i = 1,10 + ENDDO + !$acc kernels loop worker(5) + DO i = 1,10 + ENDDO + !$acc kernels loop worker(num:5) + DO i = 1,10 + ENDDO + !$acc kernels loop worker + DO i = 1,10 + !$acc kernels loop worker ! { dg-error "OpenACC construct inside of non-OpenACC region" } + DO j = 1,10 + ENDDO + !$acc kernels loop gang ! { dg-error "OpenACC construct inside of non-OpenACC region" } + DO j = 1,10 + ENDDO + ENDDO + !$acc kernels loop gang worker + DO i = 1,10 + ENDDO + + !$acc kernels loop vector + DO i = 1,10 + ENDDO + !$acc kernels loop vector(5) + DO i = 1,10 + ENDDO + !$acc kernels loop vector(length:5) + DO i = 1,10 + ENDDO + !$acc kernels loop vector + DO i = 1,10 + !$acc kernels loop vector ! { dg-error "OpenACC construct inside of non-OpenACC region" } + DO j = 1,10 + ENDDO + !$acc kernels loop worker ! { dg-error "OpenACC construct inside of non-OpenACC region" } + DO j = 1,10 + ENDDO + !$acc kernels loop gang ! { dg-error "OpenACC construct inside of non-OpenACC region" } + DO j = 1,10 + ENDDO + ENDDO + !$acc kernels loop gang vector + DO i = 1,10 + ENDDO + !$acc kernels loop worker vector + DO i = 1,10 + ENDDO + + !$acc kernels loop auto + DO i = 1,10 + ENDDO + + !$acc kernels loop tile(1) + DO i = 1,10 + ENDDO + !$acc kernels loop tile(*) + DO i = 1,10 + ENDDO + !$acc kernels loop tile(*, 1) + DO i = 1,10 + DO j = 1,10 + ENDDO + ENDDO + !$acc kernels loop tile(-1) ! { dg-warning "must be positive" } + do i = 1,10 + enddo + !$acc kernels loop vector tile(*) + DO i = 1,10 + ENDDO + !$acc kernels loop worker tile(*) + DO i = 1,10 + ENDDO + !$acc kernels loop gang tile(*) + DO i = 1,10 + ENDDO + !$acc kernels loop vector gang tile(*) + DO i = 1,10 + ENDDO + !$acc kernels loop vector worker tile(*) + DO i = 1,10 + ENDDO + !$acc kernels loop gang worker tile(*) + DO i = 1,10 + ENDDO + + !$acc parallel loop auto + DO i = 1,10 + ENDDO + !$acc parallel loop gang + DO i = 1,10 + ENDDO + !$acc parallel loop gang(static:5) + DO i = 1,10 + ENDDO + !$acc parallel loop gang(static:*) + DO i = 1,10 + ENDDO + !$acc parallel loop gang + DO i = 1,10 + !$acc parallel loop gang ! { dg-error "OpenACC construct inside of non-OpenACC region" } + DO j = 1,10 + ENDDO + ENDDO + + !$acc parallel loop worker + DO i = 1,10 + ENDDO + !$acc parallel loop worker + DO i = 1,10 + !$acc parallel loop worker ! { dg-error "OpenACC construct inside of non-OpenACC region" } + DO j = 1,10 + ENDDO + !$acc parallel loop gang ! { dg-error "OpenACC construct inside of non-OpenACC region" } + DO j = 1,10 + ENDDO + ENDDO + !$acc parallel loop gang worker + DO i = 1,10 + ENDDO + + !$acc parallel loop vector + DO i = 1,10 + !$acc parallel loop vector ! { dg-error "OpenACC construct inside of non-OpenACC region" } + DO j = 1,10 + ENDDO + !$acc parallel loop worker ! { dg-error "OpenACC construct inside of non-OpenACC region" } + DO j = 1,10 + ENDDO + !$acc parallel loop gang ! { dg-error "OpenACC construct inside of non-OpenACC region" } + DO j = 1,10 + ENDDO + ENDDO + !$acc parallel loop gang vector + DO i = 1,10 + ENDDO + !$acc parallel loop worker vector + DO i = 1,10 + ENDDO + + !$acc parallel loop auto + DO i = 1,10 + ENDDO + + !$acc parallel loop tile(1) + DO i = 1,10 + ENDDO + !$acc parallel loop tile(*) + DO i = 1,10 + ENDDO + !$acc parallel loop tile(*, 1) + DO i = 1,10 + DO j = 1,10 + ENDDO + ENDDO + !$acc parallel loop tile(-1) ! { dg-warning "must be positive" } + do i = 1,10 + enddo + !$acc parallel loop vector tile(*) + DO i = 1,10 + ENDDO + !$acc parallel loop worker tile(*) + DO i = 1,10 + ENDDO + !$acc parallel loop gang tile(*) + DO i = 1,10 + ENDDO + !$acc parallel loop vector gang tile(*) + DO i = 1,10 + ENDDO + !$acc parallel loop vector worker tile(*) + DO i = 1,10 + ENDDO + !$acc parallel loop gang worker tile(*) + DO i = 1,10 + ENDDO +end diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-6.f95 b/gcc/testsuite/gfortran.dg/goacc/loop-6.f95 new file mode 100644 index 00000000000..e13abc764c8 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/loop-6.f95 @@ -0,0 +1,82 @@ +! { dg-do compile } +! { dg-additional-options "-fmax-errors=100" } + +! This error is temporary. Remove when support is added for these clauses +! in the middle end. +! { dg-prune-output "sorry, unimplemented" } +! { dg-prune-output "Error: work-sharing region" } + +program test + implicit none + integer :: i, j + + !$acc parallel + !$acc loop auto + DO i = 1,10 + ENDDO + !$acc loop gang + DO i = 1,10 + ENDDO + !$acc loop gang(static:5) + DO i = 1,10 + ENDDO + !$acc loop gang(static:*) + DO i = 1,10 + ENDDO + !$acc loop gang + DO i = 1,10 + !$acc loop vector + DO j = 1,10 + ENDDO + !$acc loop worker + DO j = 1,10 + ENDDO + ENDDO + + !$acc loop worker + DO i = 1,10 + ENDDO + !$acc loop worker + DO i = 1,10 + !$acc loop vector + DO j = 1,10 + ENDDO + ENDDO + !$acc loop gang worker + DO i = 1,10 + ENDDO + + !$acc loop vector + DO i = 1,10 + ENDDO + !$acc loop vector(5) ! { dg-error "argument not permitted" } + DO i = 1,10 + ENDDO + !$acc loop vector(length:5) ! { dg-error "argument not permitted" } + DO i = 1,10 + ENDDO + !$acc loop vector + DO i = 1,10 + ENDDO + !$acc loop gang vector + DO i = 1,10 + ENDDO + !$acc loop worker vector + DO i = 1,10 + ENDDO + + !$acc loop auto + DO i = 1,10 + ENDDO + !$acc end parallel + + !$acc parallel loop vector + DO i = 1,10 + ENDDO + !$acc parallel loop vector(5) ! { dg-error "argument not permitted" } + DO i = 1,10 + ENDDO + !$acc parallel loop vector(length:5) ! { dg-error "argument not permitted" } + DO i = 1,10 + ENDDO +end diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90 b/gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90 index d72bae43930..6cfd715e118 100644 --- a/gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90 @@ -3,6 +3,11 @@ ! test for tree-dump-original and spaces-commas +! This error is temporary. Remove when support is added for these clauses +! in the middle end. +! { dg-prune-output "sorry, unimplemented" } +! { dg-prune-output "Error: work-sharing region" } + program test implicit none integer :: i, j, k, m, sum @@ -17,7 +22,7 @@ program test !$acc loop independent gang (3) DO i = 1,10 - !$acc loop worker(3) ! { dg-error "work-sharing region may not be closely nested inside of work-sharing, critical, ordered, master or explicit task region" } + !$acc loop worker(3) DO j = 1,10 !$acc loop vector(5) DO k = 1,10 diff --git a/gcc/testsuite/gfortran.dg/goacc/multi-clause.f90 b/gcc/testsuite/gfortran.dg/goacc/multi-clause.f90 new file mode 100644 index 00000000000..287007623c3 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/multi-clause.f90 @@ -0,0 +1,13 @@ +! Test if variable appearing in multiple clauses are errors. + +! { dg-compile } + +program combined + implicit none + integer a(100), i, j + + !$acc parallel loop reduction (+:j) copy (j) copyout(j) ! { dg-error "Symbol 'j' present on multiple clauses" } + do i = 1, 100 + end do + !$acc end parallel loop +end program combined diff --git a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 index 4915744de6b..9037f6c895f 100644 --- a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 @@ -37,4 +37,3 @@ end program test ! { dg-final { scan-tree-dump-times "map\\(force_deviceptr:u\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times "private\\(v\\)" 1 "original" } } -! { dg-final { scan-tree-dump-times "firstprivate\\(w\\)" 1 "original" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/update.f95 b/gcc/testsuite/gfortran.dg/goacc/update.f95 new file mode 100644 index 00000000000..d88d20e22d3 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/update.f95 @@ -0,0 +1,5 @@ +! { dg-do compile } + +program foo + !$acc update ! { dg-error "must contain at least one 'device' or 'host' or 'self' clause" } +end program foo diff --git a/gcc/tree-core.h b/gcc/tree-core.h index d913e510b7d..f2299f28442 100644 --- a/gcc/tree-core.h +++ b/gcc/tree-core.h @@ -432,7 +432,10 @@ enum omp_clause_code { OMP_CLAUSE_NUM_WORKERS, /* OpenACC clause: vector_length (integer-expression). */ - OMP_CLAUSE_VECTOR_LENGTH + OMP_CLAUSE_VECTOR_LENGTH, + + /* OpenACC clause: tile ( size-expr-list ). */ + OMP_CLAUSE_TILE }; #undef DEFTREESTRUCT diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c index efae4c0dd5b..b1685fd1018 100644 --- a/gcc/tree-pretty-print.c +++ b/gcc/tree-pretty-print.c @@ -932,6 +932,12 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, int flags) case OMP_CLAUSE_INDEPENDENT: pp_string (pp, "independent"); break; + case OMP_CLAUSE_TILE: + pp_string (pp, "tile("); + dump_generic_node (pp, OMP_CLAUSE_TILE_LIST (clause), + spc, flags, false); + pp_right_paren (pp); + break; default: /* Should never happen. */ diff --git a/gcc/tree.c b/gcc/tree.c index 18d6544d0b0..5b9a7bdc2fe 100644 --- a/gcc/tree.c +++ b/gcc/tree.c @@ -328,6 +328,7 @@ unsigned const char omp_clause_num_ops[] = 1, /* OMP_CLAUSE_NUM_GANGS */ 1, /* OMP_CLAUSE_NUM_WORKERS */ 1, /* OMP_CLAUSE_VECTOR_LENGTH */ + 1, /* OMP_CLAUSE_TILE */ }; const char * const omp_clause_code_name[] = @@ -398,7 +399,8 @@ const char * const omp_clause_code_name[] = "vector", "num_gangs", "num_workers", - "vector_length" + "vector_length", + "tile" }; @@ -11595,6 +11597,7 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data, case OMP_CLAUSE_DEFAULTMAP: case OMP_CLAUSE_AUTO: case OMP_CLAUSE_SEQ: + case OMP_CLAUSE_TILE: WALK_SUBTREE_TAIL (OMP_CLAUSE_CHAIN (*tp)); case OMP_CLAUSE_LASTPRIVATE: diff --git a/gcc/tree.h b/gcc/tree.h index dc592d492e2..6768b3bb29b 100644 --- a/gcc/tree.h +++ b/gcc/tree.h @@ -1556,6 +1556,9 @@ extern void protected_set_expr_location (tree, location_t); #define OMP_CLAUSE_DEFAULT_KIND(NODE) \ (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_DEFAULT)->omp_clause.subcode.default_kind) +#define OMP_CLAUSE_TILE_LIST(NODE) \ + OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_TILE), 0) + /* SSA_NAME accessors. */ /* Returns the IDENTIFIER_NODE giving the SSA name a name or NULL_TREE