[3/5] OpenACC tile clause support, C/C++ front-end parts

Message ID	61e744b3-3ae8-1709-df7b-52d65bbfb676@mentor.com
State	New
Headers	show Return-Path: <gcc-patches-return-440936-incoming=patchwork.ozlabs.org@gcc.gnu.org> DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :subject:to:cc:message-id:date:mime-version:content-type; q=dns; s=default; b=n2NYcOJJAxwBckaL3zMQGmUCFXNTdLxmC0f9PoljPl6l0sWdxA v1UNcmlewllO+ZjjEaxsB6Eh8yN5a9YF3ISD/aSPEv6ylSrU+ARSkAfITges7SBT PU6sKUx1I7J/8UVTY/adhlmUL1ZHnkl4zVvYSrs2oWJhYf1AR/SKUn75w= Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk Sender: gcc-patches-owner@gcc.gnu.org From: Chung-Lin Tang <chunglin_tang@mentor.com> Subject: [Patch 3/5] OpenACC tile clause support, C/C++ front-end parts To: gcc-patches <gcc-patches@gcc.gnu.org>, Jakub Jelinek <jakub@redhat.com> CC: Nathan Sidwell <nathan@acm.org>, Cesar Philippidis <cesar@codesourcery.com>, Chung-Lin Tang <chunglin_tang@mentor.com> Message-ID: <61e744b3-3ae8-1709-df7b-52d65bbfb676@mentor.com> Date: Thu, 10 Nov 2016 18:46:16 +0800 User-Agent: Mozilla/5.0 (Macintosh; Intel Mac OS X 10.11; rv:45.0) Gecko/20100101 Thunderbird/45.4.0 MIME-Version: 1.0 Content-Type: multipart/mixed; boundary="------------EDAF4EB0DC7037A4A0D598E1"

Index: c/c-parser.c =================================================================== --- c/c-parser.c (revision 241809) +++ c/c-parser.c (working copy) @@ -11010,6 +11010,7 @@ c_parser_omp_clause_collapse (c_parser *parser, tr location_t loc; check_no_duplicate_clause (list, OMP_CLAUSE_COLLAPSE, "collapse"); + 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 %<(%>")) @@ -11920,10 +11921,11 @@ static tree c_parser_oacc_clause_tile (c_parser *parser, tree list) { tree c, expr = error_mark_node; - location_t loc, expr_loc; + location_t loc; tree tile = NULL_TREE; check_no_duplicate_clause (list, OMP_CLAUSE_TILE, "tile"); + check_no_duplicate_clause (list, OMP_CLAUSE_COLLAPSE, "collapse"); loc = c_parser_peek_token (parser)->location; if (!c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>")) @@ -11931,16 +11933,19 @@ c_parser_oacc_clause_tile (c_parser *parser, tree do { + if (tile && !c_parser_require (parser, CPP_COMMA, "expected %<,%>")) + return list; + 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; + expr = integer_zero_node; } else { - expr_loc = c_parser_peek_token (parser)->location; + location_t expr_loc = c_parser_peek_token (parser)->location; c_expr cexpr = c_parser_expr_no_commas (parser, NULL); cexpr = convert_lvalue_to_rvalue (expr_loc, cexpr, false, true); expr = cexpr.value; @@ -11952,28 +11957,20 @@ c_parser_oacc_clause_tile (c_parser *parser, tree return list; } - if (!INTEGRAL_TYPE_P (TREE_TYPE (expr))) - { - c_parser_error (parser, "%<tile%> value must be integral"); - return list; - } - 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) + if (!INTEGRAL_TYPE_P (TREE_TYPE (expr)) + || TREE_CODE (expr) != INTEGER_CST + || !tree_fits_shwi_p (expr) + || tree_to_shwi (expr) <= 0) { - warning_at (expr_loc, 0,"%<tile%> value must be positive"); - expr = integer_one_node; + error_at (expr_loc, "%<tile%> argument needs positive" + " integral constant"); + expr = integer_zero_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)); @@ -14899,11 +14896,17 @@ c_parser_omp_for_loop (location_t loc, c_parser *p bool fail = false, open_brace_parsed = false; int i, collapse = 1, ordered = 0, count, nbraces = 0; location_t for_loc; + bool tiling = false; vec<tree, va_gc> *for_block = make_tree_vector (); for (cl = clauses; cl; cl = OMP_CLAUSE_CHAIN (cl)) if (OMP_CLAUSE_CODE (cl) == OMP_CLAUSE_COLLAPSE) collapse = tree_to_shwi (OMP_CLAUSE_COLLAPSE_EXPR (cl)); + else if (OMP_CLAUSE_CODE (cl) == OMP_CLAUSE_TILE) + { + tiling = true; + collapse = list_length (OMP_CLAUSE_TILE_LIST (cl)); + } else if (OMP_CLAUSE_CODE (cl) == OMP_CLAUSE_ORDERED && OMP_CLAUSE_ORDERED_EXPR (cl)) { @@ -14933,7 +14936,7 @@ c_parser_omp_for_loop (location_t loc, c_parser *p pc = &OMP_CLAUSE_CHAIN (*pc); } - gcc_assert (collapse >= 1 && ordered >= 0); + gcc_assert (tiling || (collapse >= 1 && ordered >= 0)); count = ordered ? ordered : collapse; declv = make_tree_vec (count); Index: cp/pt.c =================================================================== --- cp/pt.c (revision 241809) +++ cp/pt.c (working copy) @@ -14713,6 +14713,7 @@ tsubst_omp_clauses (tree clauses, enum c_omp_regio nc = copy_node (oc); OMP_CLAUSE_CHAIN (nc) = new_clauses; new_clauses = nc; + bool needs_ice = false; switch (OMP_CLAUSE_CODE (nc)) { @@ -14742,10 +14743,16 @@ tsubst_omp_clauses (tree clauses, enum c_omp_regio = tsubst_omp_clause_decl (OMP_CLAUSE_DECL (oc), args, complain, in_decl); break; + case OMP_CLAUSE_COLLAPSE: + case OMP_CLAUSE_TILE: + /* These clauses really need a positive integral constant + expression, but we don't have a predicate for that + (yet). */ + needs_ice = true; + /* FALLTHRU */ case OMP_CLAUSE_IF: case OMP_CLAUSE_NUM_THREADS: case OMP_CLAUSE_SCHEDULE: - case OMP_CLAUSE_COLLAPSE: case OMP_CLAUSE_FINAL: case OMP_CLAUSE_DEVICE: case OMP_CLAUSE_DIST_SCHEDULE: @@ -14766,8 +14773,8 @@ tsubst_omp_clauses (tree clauses, enum c_omp_regio 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); + = tsubst_expr (OMP_CLAUSE_OPERAND (oc, 0), args, complain, in_decl, + /*integral_constant_expression_p=*/needs_ice); break; case OMP_CLAUSE_REDUCTION: if (OMP_CLAUSE_REDUCTION_PLACEHOLDER (oc)) @@ -14836,19 +14843,6 @@ tsubst_omp_clauses (tree clauses, enum c_omp_regio 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 (); } Index: cp/semantics.c =================================================================== --- cp/semantics.c (revision 241809) +++ cp/semantics.c (working copy) @@ -7086,7 +7086,8 @@ finish_omp_clauses (tree clauses, enum c_omp_regio else if (!type_dependent_expression_p (t) && !INTEGRAL_TYPE_P (TREE_TYPE (t))) { - error ("%<tile%> value must be integral"); + error_at (OMP_CLAUSE_LOCATION (c), + "%<tile%> argument needs integral type"); remove = true; } else @@ -7094,14 +7095,16 @@ finish_omp_clauses (tree clauses, enum c_omp_regio t = mark_rvalue_use (t); if (!processing_template_decl) { + /* Zero is used to indicate '*', we permit you + to get there via an ICE of value zero. */ t = maybe_constant_value (t); - if (TREE_CODE (t) == INTEGER_CST - && tree_int_cst_sgn (t) != 1 - && t != integer_minus_one_node) + if (TREE_CODE (t) != INTEGER_CST + || !tree_fits_shwi_p (t) + || tree_to_shwi (t) < 0) { - warning_at (OMP_CLAUSE_LOCATION (c), 0, - "%<tile%> value must be positive"); - t = integer_one_node; + error_at (OMP_CLAUSE_LOCATION (c), + "%<tile%> argument needs positive integral constant"); + remove = true; } } t = fold_build_cleanup_point_expr (TREE_TYPE (t), t); @@ -8000,11 +8003,19 @@ finish_omp_for (location_t locus, enum tree_code c gcc_assert (TREE_VEC_LENGTH (declv) == TREE_VEC_LENGTH (incrv)); if (TREE_VEC_LENGTH (declv) > 1) { - tree c = find_omp_clause (clauses, OMP_CLAUSE_COLLAPSE); + tree c; + + c = find_omp_clause (clauses, OMP_CLAUSE_TILE); if (c) - collapse = tree_to_shwi (OMP_CLAUSE_COLLAPSE_EXPR (c)); - if (collapse != TREE_VEC_LENGTH (declv)) - ordered = TREE_VEC_LENGTH (declv); + collapse = list_length (OMP_CLAUSE_TILE_LIST (c)); + else + { + c = find_omp_clause (clauses, OMP_CLAUSE_COLLAPSE); + if (c) + collapse = tree_to_shwi (OMP_CLAUSE_COLLAPSE_EXPR (c)); + if (collapse != TREE_VEC_LENGTH (declv)) + ordered = TREE_VEC_LENGTH (declv); + } } for (i = 0; i < TREE_VEC_LENGTH (declv); i++) { Index: cp/parser.c =================================================================== --- cp/parser.c (revision 241809) +++ cp/parser.c (working copy) @@ -30885,30 +30885,33 @@ cp_parser_oacc_clause_tile (cp_parser *parser, loc tree c, expr = error_mark_node; tree tile = NULL_TREE; + /* Collapse and tile are mutually exclusive. (The spec doesn't say + so, but the spec authors never considered such a case and have + differing opinions on what it might mean, including 'not + allowed'.) */ check_no_duplicate_clause (list, OMP_CLAUSE_TILE, "tile", clause_loc); + check_no_duplicate_clause (list, OMP_CLAUSE_COLLAPSE, "collapse", + clause_loc); if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN)) return list; do { + if (tile && !cp_parser_require (parser, CPP_COMMA, RT_COMMA)) + return list; + 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; + expr = integer_zero_node; } else - expr = cp_parser_assignment_expression (parser, NULL, false, false); + expr = cp_parser_constant_expression (parser); - 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)); @@ -31021,6 +31024,7 @@ cp_parser_omp_clause_collapse (cp_parser *parser, } check_no_duplicate_clause (list, OMP_CLAUSE_COLLAPSE, "collapse", location); + check_no_duplicate_clause (list, OMP_CLAUSE_TILE, "tile", location); c = build_omp_clause (loc, OMP_CLAUSE_COLLAPSE); OMP_CLAUSE_CHAIN (c) = list; OMP_CLAUSE_COLLAPSE_EXPR (c) = num; @@ -34027,10 +34031,16 @@ cp_parser_omp_for_loop (cp_parser *parser, enum tr int i, collapse = 1, ordered = 0, count, nbraces = 0; vec<tree, va_gc> *for_block = make_tree_vector (); auto_vec<tree, 4> orig_inits; + bool tiling = false; for (cl = clauses; cl; cl = OMP_CLAUSE_CHAIN (cl)) if (OMP_CLAUSE_CODE (cl) == OMP_CLAUSE_COLLAPSE) collapse = tree_to_shwi (OMP_CLAUSE_COLLAPSE_EXPR (cl)); + else if (OMP_CLAUSE_CODE (cl) == OMP_CLAUSE_TILE) + { + tiling = true; + collapse = list_length (OMP_CLAUSE_TILE_LIST (cl)); + } else if (OMP_CLAUSE_CODE (cl) == OMP_CLAUSE_ORDERED && OMP_CLAUSE_ORDERED_EXPR (cl)) { @@ -34060,7 +34070,7 @@ cp_parser_omp_for_loop (cp_parser *parser, enum tr pc = &OMP_CLAUSE_CHAIN (*pc); } - gcc_assert (collapse >= 1 && ordered >= 0); + gcc_assert (tiling || (collapse >= 1 && ordered >= 0)); count = ordered ? ordered : collapse; declv = make_tree_vec (count); @@ -34079,13 +34089,15 @@ cp_parser_omp_for_loop (cp_parser *parser, enum tr if (code != CILK_FOR && !cp_lexer_next_token_is_keyword (parser->lexer, RID_FOR)) { - cp_parser_error (parser, "for statement expected"); + if (!collapse_err) + cp_parser_error (parser, "for statement expected"); return NULL; } if (code == CILK_FOR && !cp_lexer_next_token_is_keyword (parser->lexer, RID_CILK_FOR)) { - cp_parser_error (parser, "_Cilk_for statement expected"); + if (!collapse_err) + cp_parser_error (parser, "_Cilk_for statement expected"); return NULL; } loc = cp_lexer_consume_token (parser->lexer)->location; @@ -34245,7 +34257,7 @@ cp_parser_omp_for_loop (cp_parser *parser, enum tr nested. Hopefully the final version clarifies this. For now handle (multiple) {'s and empty statements. */ cp_parser_parse_tentatively (parser); - do + for (;;) { if (cp_lexer_next_token_is_keyword (parser->lexer, RID_FOR)) break; @@ -34260,14 +34272,13 @@ cp_parser_omp_for_loop (cp_parser *parser, enum tr else { loc = cp_lexer_peek_token (parser->lexer)->location; - error_at (loc, "not enough collapsed for loops"); + error_at (loc, "not enough for loops to collapse"); collapse_err = true; cp_parser_abort_tentative_parse (parser); declv = NULL_TREE; break; } } - while (1); if (declv) { Index: testsuite/c-c++-common/goacc/loop-auto-2.c =================================================================== --- testsuite/c-c++-common/goacc/loop-auto-2.c (revision 0) +++ testsuite/c-c++-common/goacc/loop-auto-2.c (revision 0) @@ -0,0 +1,107 @@ + +// Tile parititioning + +void Ok () +{ +#pragma acc parallel num_gangs (10) num_workers(32) vector_length(32) + { + +#pragma acc loop tile(*) gang vector + for (int ix = 0; ix < 10; ix++) + { + } + +#pragma acc loop tile(*) + for (int ix = 0; ix < 10; ix++) + { + } + +#pragma acc loop tile(*) gang + for (int ix = 0; ix < 10; ix++) + { + #pragma acc loop vector + for (int jx = 0; jx < 10; jx++) + ; + } + +#pragma acc loop tile(*) + for (int ix = 0; ix < 10; ix++) + { + #pragma acc loop vector + for (int jx = 0; jx < 10; jx++) + ; + } + +#pragma acc loop gang + for (int jx = 0; jx < 10; jx++) + { +#pragma acc loop tile(*) vector + for (int ix = 0; ix < 10; ix++) + { + } + +#pragma acc loop tile(*) + for (int ix = 0; ix < 10; ix++) + { + } + } + +#pragma acc loop tile(*) worker + for (int ix = 0; ix < 10; ix++) + { + #pragma acc loop vector + for (int jx = 0; jx < 10; jx++) + ; + } + } +} + +void Bad () +{ +#pragma acc parallel num_gangs (10) num_workers(32) vector_length(32) + { + +#pragma acc loop tile(*) gang vector /* { dg-message "containing loop" } */ + for (int ix = 0; ix < 10; ix++) + { +#pragma acc loop vector /* { dg-error "uses same" } */ + for (int jx = 0; jx < 10; jx++) + ; + } + +#pragma acc loop tile(*) gang vector + for (int ix = 0; ix < 10; ix++) + { + #pragma acc loop auto /* { dg-warning "insufficient partitioning" } */ + for (int jx = 0; jx < 10; jx++) + ; + } + +#pragma acc loop tile(*) auto /* { dg-warning "insufficient partitioning" } */ + for (int ix = 0; ix < 10; ix++) + { + #pragma acc loop worker + for (int jx = 0; jx < 10; jx++) + ; + } + +#pragma acc loop worker /* { dg-message "containing loop" } */ + for (int jx = 0; jx < 10; jx++) + { +#pragma acc loop tile(*) gang vector /* { dg-error "incorrectly nested" } */ + for (int ix = 0; ix < 10; ix++) + { + } + +#pragma acc loop tile(*) vector /* { dg-warning "insufficient partitioning" } */ + for (int ix = 0; ix < 10; ix++) + { + } + +#pragma acc loop tile(*) /* { dg-warning "insufficient partitioning" } */ + for (int ix = 0; ix < 10; ix++) + { + } + } + } +} Index: testsuite/c-c++-common/goacc/tile-2.c =================================================================== --- testsuite/c-c++-common/goacc/tile-2.c (revision 0) +++ testsuite/c-c++-common/goacc/tile-2.c (revision 0) @@ -0,0 +1,21 @@ +int main () +{ +#pragma acc parallel + { +#pragma acc loop tile (*,*) + for (int ix = 0; ix < 30; ix++) + ; /* { dg-error "not enough" } */ + +#pragma acc loop tile (*,*) + for (int ix = 0; ix < 30; ix++) + for (int jx = 0; jx < ix; jx++) /* { dg-error "condition expression" } */ + ; + +#pragma acc loop tile (*) + for (int ix = 0; ix < 30; ix++) + for (int jx = 0; jx < ix; jx++) /* OK */ + ; + + } + return 0; +} Index: testsuite/c-c++-common/goacc/loop-auto-1.c =================================================================== --- testsuite/c-c++-common/goacc/loop-auto-1.c (revision 241809) +++ testsuite/c-c++-common/goacc/loop-auto-1.c (working copy) @@ -74,6 +74,21 @@ void Foo () for (int kx = 0; kx < 10; kx++) {} } } + +#pragma acc loop auto + for (int ix = 0; ix < 10; ix++) + { +#pragma acc loop auto + for (int jx = 0; jx < 10; jx++) + { +#pragma acc loop auto /* { dg-warning "insufficient partitioning" } */ + for (int kx = 0; kx < 10; kx++) + { +#pragma acc loop auto + for (int lx = 0; lx < 10; lx++) {} + } + } + } } } @@ -214,10 +229,10 @@ void Vector (void) #pragma acc loop auto for (int ix = 0; ix < 10; ix++) {} -#pragma acc loop auto +#pragma acc loop auto /* { dg-warning "insufficient partitioning" } */ for (int ix = 0; ix < 10; ix++) { -#pragma acc loop auto /* { dg-warning "insufficient partitioning" } */ +#pragma acc loop auto for (int jx = 0; jx < 10; jx++) {} } } Index: testsuite/c-c++-common/goacc/tile.c =================================================================== --- testsuite/c-c++-common/goacc/tile.c (revision 241809) +++ testsuite/c-c++-common/goacc/tile.c (working copy) @@ -1,7 +1,9 @@ +#include <stdbool.h> + int main () { - int i, *a, b; + int i, j, k, *a, b; #pragma acc parallel loop tile (10) for (i = 0; i < 100; i++) @@ -13,11 +15,14 @@ main () #pragma acc parallel loop tile (10, *) for (i = 0; i < 100; i++) - ; + for (j = 0; j < 100; j++) + ; -#pragma acc parallel loop tile (10, *, i) +#pragma acc parallel loop tile (10, *, i) // { dg-error "" } for (i = 0; i < 100; i++) - ; + for (j = 0; j < 100; j++) + for (k = 0; k < 100; k++) + ; #pragma acc parallel loop tile // { dg-error "expected '\\\('" } for (i = 0; i < 100; i++) @@ -35,37 +40,44 @@ main () for (i = 0; i < 100; i++) ; -#pragma acc parallel loop tile (1.1) // { dg-error "'tile' value must be integral" } +#pragma acc parallel loop tile (1.1) // { dg-error "'tile' argument needs" } for (i = 0; i < 100; i++) ; -#pragma acc parallel loop tile (-3) // { dg-warning "'tile' value must be positive" } +#pragma acc parallel loop tile (-3) // { dg-error "'tile' argument needs" } for (i = 0; i < 100; i++) ; -#pragma acc parallel loop tile (10,-3) // { dg-warning "'tile' value must be positive" } +#pragma acc parallel loop tile (10,-3) // { dg-error "'tile' argument needs" } for (i = 0; i < 100; i++) - ; + for (j = 0; j < 100; j++) + ; -#pragma acc parallel loop tile (-100,10,5) // { dg-warning "'tile' value must be positive" } +#pragma acc parallel loop tile (-100,10,5) // { dg-error "'tile' argument needs" } for (i = 0; i < 100; i++) - ; + for (j = 0; j < 100; j++) + for (k = 0; k < 100; k++) + ; -#pragma acc parallel loop tile (1,2.0,true) // { dg-error "" } +#pragma acc parallel loop tile (1,true) for (i = 0; i < 100; i++) - ; + for (j = 0; j < 100; j++) + ; -#pragma acc parallel loop tile (*a, 1) +#pragma acc parallel loop tile (*a, 1) // { dg-error "" } for (i = 0; i < 100; i++) - ; + for (j = 0; j < 100; j++) + ; -#pragma acc parallel loop tile (1, *a, b) +#pragma acc parallel loop tile (1, b) // { dg-error "" } for (i = 0; i < 100; i++) - ; + for (j = 0; j < 100; j++) + ; -#pragma acc parallel loop tile (b, 1, *a) +#pragma acc parallel loop tile (b, 1) // { dg-error "" } for (i = 0; i < 100; i++) - ; + for (j = 0; j < 100; j++) + ; return 0; } @@ -73,7 +85,7 @@ main () void par (void) { - int i, j; + int i, j, k; #pragma acc parallel { @@ -95,22 +107,22 @@ void par (void) for (j = 1; j < 10; j++) { } } -#pragma acc loop tile(-2) // { dg-warning "'tile' value must be positive" } +#pragma acc loop tile(-2) // { dg-error "'tile' argument needs" } for (i = 1; i < 10; i++) { } -#pragma acc loop tile(i) +#pragma acc loop tile(i) // { dg-error "" } for (i = 1; i < 10; i++) { } #pragma acc loop tile(2, 2, 1) for (i = 1; i < 3; i++) { for (j = 4; j < 6; j++) - { } + for (k = 0; k< 100; k++); } #pragma acc loop tile(2, 2) for (i = 1; i < 5; i+=2) { - for (j = i + 1; j < 7; j+=i) + for (j = i + 1; j < 7; j+=i) // { dg-error "initializer expression" } { } } #pragma acc loop vector tile(*) @@ -156,24 +168,21 @@ void p3 (void) for (j = 1; j < 10; j++) { } } -#pragma acc parallel loop tile(-2) // { dg-warning "'tile' value must be positive" } +#pragma acc parallel loop tile(-2) // { dg-error "'tile' argument needs" } for (i = 1; i < 10; i++) { } -#pragma acc parallel loop tile(i) +#pragma acc parallel loop tile(i) // { dg-error "" } for (i = 1; i < 10; i++) { } #pragma acc parallel loop tile(2, 2, 1) for (i = 1; i < 3; i++) - { - for (j = 4; j < 6; j++) - { } - } + for (j = 4; j < 6; j++) + for (int k = 1 ; k < 2; k++) + ; #pragma acc parallel loop tile(2, 2) for (i = 1; i < 5; i+=2) - { - for (j = i + 1; j < 7; j++) - { } - } + for (j = i + 1; j < 7; j++) // { dg-error "initializer expression" } + { } #pragma acc parallel loop vector tile(*) for (i = 0; i < 10; i++) { } @@ -227,22 +236,23 @@ kern (void) #pragma acc loop tile(*, 1) for (i = 0; i < 10; i++) { - for (j = 0; j < 10; i++) + for (j = 0; j < 10; i++) /* { dg-error "increment expression" } */ { } } -#pragma acc loop tile(-2) // { dg-warning "'tile' value must be positive" } +#pragma acc loop tile(-2) // { dg-error "'tile' argument needs" } for (i = 0; i < 10; i++) { } -#pragma acc loop tile(i) +#pragma acc loop tile(i) // { dg-error "" } for (i = 0; i < 10; i++) { } #pragma acc loop tile(2, 2, 1) for (i = 2; i < 4; i++) - for (i = 4; i < 6; i++) + for (j = 4; j < 6; j++) + for (int k = 4; k < 6; k++) { } #pragma acc loop tile(2, 2) for (i = 1; i < 5; i+=2) - for (j = i+1; j < 7; i++) + for (j = i+1; j < 7; j++) /* { dg-error "initializer expression" } */ { } #pragma acc loop vector tile(*) for (i = 0; i < 10; i++) @@ -288,22 +298,21 @@ void k3 (void) for (j = 1; j < 10; j++) { } } -#pragma acc kernels loop tile(-2) // { dg-warning "'tile' value must be positive" } +#pragma acc kernels loop tile(-2) // { dg-error "'tile' argument needs" } for (i = 1; i < 10; i++) { } -#pragma acc kernels loop tile(i) +#pragma acc kernels loop tile(i) // { dg-error "" } for (i = 1; i < 10; i++) { } #pragma acc kernels loop tile(2, 2, 1) for (i = 1; i < 3; i++) - { - for (j = 4; j < 6; j++) - { } - } + for (j = 4; j < 6; j++) + for (int k = 1; k < 7; k++) + ; #pragma acc kernels loop tile(2, 2) for (i = 1; i < 5; i++) { - for (j = i + 1; j < 7; j += i) + for (j = i + 1; j < 7; j += i) /* { dg-error "initializer expression" } */ { } } #pragma acc kernels loop vector tile(*) Index: testsuite/g++.dg/goacc/template.C =================================================================== --- testsuite/g++.dg/goacc/template.C (revision 241809) +++ testsuite/g++.dg/goacc/template.C (working copy) @@ -5,7 +5,7 @@ accDouble(int val) return val * 2; } -template<typename T> T +template<typename T, int I> T oacc_parallel_copy (T a) { T b = 0; @@ -36,7 +36,7 @@ oacc_parallel_copy (T a) for (int j = 0; j < 5; j++) b = a; -#pragma acc loop auto tile (a, 3) +#pragma acc loop auto tile (I, 3) for (int i = 0; i < a; i++) for (int j = 0; j < 5; j++) b = a; @@ -135,7 +135,7 @@ oacc_kernels_copy (T a) int main () { - int b = oacc_parallel_copy<int> (5); + int b = oacc_parallel_copy<int, 4> (5); int c = oacc_kernels_copy<int> (5); return b + c; Index: testsuite/g++.dg/goacc/tile-1.C =================================================================== --- testsuite/g++.dg/goacc/tile-1.C (revision 0) +++ testsuite/g++.dg/goacc/tile-1.C (revision 0) @@ -0,0 +1,16 @@ +/* of tile erroneously clobbered the template, resulting + in missing errors and other fun. */ + +template <int I> +void Foo () +{ +#pragma acc parallel loop tile(I) // { dg-error "" } + for (int ix = 0; ix < 10; ix++) + ; +} + +int main () +{ + Foo<1> (); // OK + Foo<-1> (); // error +} Index: testsuite/gcc.dg/goacc/loop-processing-1.c =================================================================== --- testsuite/gcc.dg/goacc/loop-processing-1.c (revision 241809) +++ testsuite/gcc.dg/goacc/loop-processing-1.c (working copy) @@ -15,4 +15,4 @@ void vector_1 (int *ary, int size) } } -/* { dg-final { scan-tree-dump {OpenACC loops.*Loop 0\(0\).*Loop 14\(1\).*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_HEAD_MARK, 0, 1, 20\);.*Head-0:.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_HEAD_MARK, 0, 1, 20\);.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_FORK, \.data_dep\.[0-9_]+, 0\);.*Tail-0:.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_TAIL_MARK, \.data_dep\.[0-9_]+, 1\);.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_JOIN, \.data_dep\.[0-9_]+, 0\);.*Loop 6\(4\).*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_HEAD_MARK, 0, 1, 6\);.*Head-0:.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_HEAD_MARK, 0, 1, 6\);.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_FORK, \.data_dep\.[0-9_]+, 2\);.*Tail-0:.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_TAIL_MARK, \.data_dep\.[0-9_]+, 1\);.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_JOIN, \.data_dep\.[0-9_]+, 2\);} "oaccdevlow" } } */ +/* { dg-final { scan-tree-dump {OpenACC loops.*Loop 0\(0\).*Loop 24\(1\).*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_HEAD_MARK, 0, 1, 36\);.*Head-0:.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_HEAD_MARK, 0, 1, 36\);.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_FORK, \.data_dep\.[0-9_]+, 0\);.*Tail-0:.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_TAIL_MARK, \.data_dep\.[0-9_]+, 1\);.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_JOIN, \.data_dep\.[0-9_]+, 0\);.*Loop 6\(4\).*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_HEAD_MARK, 0, 2, 4\);.*Head-0:.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_HEAD_MARK, 0, 2, 4\);.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_FORK, \.data_dep\.[0-9_]+, 1\);.*Head-1:.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_HEAD_MARK, \.data_dep\.[0-9_]+, 1\);.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_FORK, \.data_dep\.[0-9_]+, 2\);.*Tail-1:.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_TAIL_MARK, \.data_dep\.[0-9_]+, 2\);.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_JOIN, \.data_dep\.[0-9_]+, 2\);.*Tail-0:.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_TAIL_MARK, \.data_dep\.[0-9_]+, 1\);.*\.data_dep\.[0-9_]+ = UNIQUE \(OACC_JOIN, \.data_dep\.[0-9_]+, 1\);} "oaccdevlow" } } */

[3/5] OpenACC tile clause support, C/C++ front-end parts

Commit Message

Comments

Patch