From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 1534) id A4417385E447; Fri, 14 May 2021 08:51:19 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org A4417385E447 Content-Type: text/plain; charset="us-ascii" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit From: Tobias Burnus To: gcc-cvs@gcc.gnu.org Subject: [gcc/devel/omp/gcc-11] OpenMP: Add support for 'close' in map clause X-Act-Checkin: gcc X-Git-Author: Marcel Vollweiler X-Git-Refname: refs/heads/devel/omp/gcc-11 X-Git-Oldrev: 5f654d3d12446c50cd6db54b75ea8966b57651bf X-Git-Newrev: c24342b9bd52cfba170a5a5e105a780db7ed5d29 Message-Id: <20210514085119.A4417385E447@sourceware.org> Date: Fri, 14 May 2021 08:51:19 +0000 (GMT) X-BeenThere: gcc-cvs@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-cvs mailing list List-Unsubscribe: , List-Archive: List-Help: List-Subscribe: , X-List-Received-Date: Fri, 14 May 2021 08:51:19 -0000 https://gcc.gnu.org/g:c24342b9bd52cfba170a5a5e105a780db7ed5d29 commit c24342b9bd52cfba170a5a5e105a780db7ed5d29 Author: Marcel Vollweiler Date: Fri May 14 10:15:38 2021 +0200 OpenMP: Add support for 'close' in map clause gcc/c/ChangeLog: * c-parser.c (c_parser_omp_clause_map): Support map-type-modifier 'close'. gcc/cp/ChangeLog: * parser.c (cp_parser_omp_clause_map): Support map-type-modifier 'close'. gcc/testsuite/ChangeLog: * c-c++-common/gomp/map-6.c: New test. * c-c++-common/gomp/map-7.c: New test. (cherry picked from commit fa6894ec9ce25f5aff10ec176212383f5c88b1ec) Diff: --- gcc/ChangeLog.omp | 16 ++++ gcc/c/ChangeLog.omp | 32 +++++--- gcc/c/c-parser.c | 116 +++++++++++++-------------- gcc/cp/ChangeLog.omp | 8 ++ gcc/cp/parser.c | 90 ++++++++++++++++----- gcc/testsuite/ChangeLog.omp | 8 ++ gcc/testsuite/c-c++-common/gomp/map-6.c | 135 ++++++++++++++++++++++++++++++++ gcc/testsuite/c-c++-common/gomp/map-7.c | 20 +++++ 8 files changed, 335 insertions(+), 90 deletions(-) diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp index 5fedea5a7ae..799a018c61b 100644 --- a/gcc/ChangeLog.omp +++ b/gcc/ChangeLog.omp @@ -1,7 +1,23 @@ 2021-05-14 Tobias Burnus Backported from master: + 2021-05-12 Tom de Vries + + PR target/96005 + * config/nvptx/nvptx-opts.h (enum ptx_version): New enum. + * config/nvptx/nvptx.c (nvptx_file_start): Print .version according + to ptx_version_option. + * config/nvptx/nvptx.h (TARGET_PTX_6_3): Define. + * config/nvptx/nvptx.md (define_insn "nvptx_shuffle") + (define_insn "nvptx_vote_ballot"): Use sync variant for + TARGET_PTX_6_3. + * config/nvptx/nvptx.opt (ptx_version): Add enum. + (mptx): Add option. + * doc/invoke.texi (Nvidia PTX Options): Add mptx item. +2021-05-14 Tobias Burnus + + Backported from master: 2021-05-07 Tobias Burnus Tom de Vries diff --git a/gcc/c/ChangeLog.omp b/gcc/c/ChangeLog.omp index ed366d518c4..f61d83aece7 100644 --- a/gcc/c/ChangeLog.omp +++ b/gcc/c/ChangeLog.omp @@ -1,3 +1,11 @@ +2021-05-14 Tobias Burnus + + Backported from master: + 2021-05-12 Marcel Vollweiler + + * c-parser.c (c_parser_omp_clause_map): Support map-type-modifier + 'close'. + 2021-05-14 Tobias Burnus Backported from master: @@ -25,22 +33,22 @@ nvptx_expand_to_rtl_hook. 2018-12-13 Cesar Philippidis - Nathan Sidwell - Julian Brown + Nathan Sidwell + Julian Brown - * c-parser.c (c_parser_omp_variable_list): New c_omp_region_type - argument. Use it to specialize handling of OMP_CLAUSE_REDUCTION for - OpenACC. + * c-parser.c (c_parser_omp_variable_list): New c_omp_region_type + argument. Use it to specialize handling of OMP_CLAUSE_REDUCTION for + OpenACC. (c_parser_oacc_data_clause): Add region-type argument. (c_parser_oacc_data_clause_deviceptr): Likewise. - (c_parser_omp_clause_reduction): Change is_omp boolean parameter to - c_omp_region_type. Update call to c_parser_omp_variable_list. - (c_parser_oacc_all_clauses): Update calls to - c_parser_omp_clause_reduction. - (c_parser_omp_all_clauses): Likewise. + (c_parser_omp_clause_reduction): Change is_omp boolean parameter to + c_omp_region_type. Update call to c_parser_omp_variable_list. + (c_parser_oacc_all_clauses): Update calls to + c_parser_omp_clause_reduction. + (c_parser_omp_all_clauses): Likewise. (c_parser_oacc_cache): Update call to c_parser_omp_var_list_parens. - * c-typeck.c (c_finish_omp_clauses): Emit an error on orphan OpenACC - gang reductions. Suppress user-defined reduction error for OpenACC. + * c-typeck.c (c_finish_omp_clauses): Emit an error on orphan OpenACC + gang reductions. Suppress user-defined reduction error for OpenACC. 2018-10-02 Thomas Schwinge Cesar Philippidis diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index 8c91822ef1e..944d503bcfc 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -15659,54 +15659,83 @@ c_parser_omp_clause_depend (c_parser *parser, tree list) map-kind: alloc | to | from | tofrom | release | delete - map ( always [,] map-kind: variable-list ) */ + map ( always [,] map-kind: variable-list ) + + OpenMP 5.0: + map ( [map-type-modifier[,] ...] map-kind: variable-list ) + + map-type-modifier: + always | close */ static tree c_parser_omp_clause_map (c_parser *parser, tree list) { location_t clause_loc = c_parser_peek_token (parser)->location; enum gomp_map_kind kind = GOMP_MAP_TOFROM; - int always = 0; - enum c_id_kind always_id_kind = C_ID_NONE; - location_t always_loc = UNKNOWN_LOCATION; - tree always_id = NULL_TREE; tree nl, c; matching_parens parens; if (!parens.require_open (parser)) return list; - if (c_parser_next_token_is (parser, CPP_NAME)) + int pos = 1; + int map_kind_pos = 0; + while (c_parser_peek_nth_token_raw (parser, pos)->type == CPP_NAME) + { + if (c_parser_peek_nth_token_raw (parser, pos + 1)->type == CPP_COLON) + { + map_kind_pos = pos; + break; + } + + if (c_parser_peek_nth_token_raw (parser, pos + 1)->type == CPP_COMMA) + pos++; + pos++; + } + + int always_modifier = 0; + int close_modifier = 0; + for (int pos = 1; pos < map_kind_pos; ++pos) { c_token *tok = c_parser_peek_token (parser); + + if (tok->type == CPP_COMMA) + { + c_parser_consume_token (parser); + continue; + } + const char *p = IDENTIFIER_POINTER (tok->value); - always_id_kind = tok->id_kind; - always_loc = tok->location; - always_id = tok->value; if (strcmp ("always", p) == 0) { - c_token *sectok = c_parser_peek_2nd_token (parser); - if (sectok->type == CPP_COMMA) + if (always_modifier) { - c_parser_consume_token (parser); - c_parser_consume_token (parser); - always = 2; + c_parser_error (parser, "too many % modifiers"); + parens.skip_until_found_close (parser); + return list; } - else if (sectok->type == CPP_NAME) + always_modifier++; + } + else if (strcmp ("close", p) == 0) + { + if (close_modifier) { - p = IDENTIFIER_POINTER (sectok->value); - if (strcmp ("alloc", p) == 0 - || strcmp ("to", p) == 0 - || strcmp ("from", p) == 0 - || strcmp ("tofrom", p) == 0 - || strcmp ("release", p) == 0 - || strcmp ("delete", p) == 0) - { - c_parser_consume_token (parser); - always = 1; - } + c_parser_error (parser, "too many % modifiers"); + parens.skip_until_found_close (parser); + return list; } + close_modifier++; } + else + { + c_parser_error (parser, "%<#pragma omp target%> with " + "modifier other than % or %" + "on % clause"); + parens.skip_until_found_close (parser); + return list; + } + + c_parser_consume_token (parser); } if (c_parser_next_token_is (parser, CPP_NAME) @@ -15716,11 +15745,11 @@ c_parser_omp_clause_map (c_parser *parser, tree list) if (strcmp ("alloc", p) == 0) kind = GOMP_MAP_ALLOC; else if (strcmp ("to", p) == 0) - kind = always ? GOMP_MAP_ALWAYS_TO : GOMP_MAP_TO; + kind = always_modifier ? GOMP_MAP_ALWAYS_TO : GOMP_MAP_TO; else if (strcmp ("from", p) == 0) - kind = always ? GOMP_MAP_ALWAYS_FROM : GOMP_MAP_FROM; + kind = always_modifier ? GOMP_MAP_ALWAYS_FROM : GOMP_MAP_FROM; else if (strcmp ("tofrom", p) == 0) - kind = always ? GOMP_MAP_ALWAYS_TOFROM : GOMP_MAP_TOFROM; + kind = always_modifier ? GOMP_MAP_ALWAYS_TOFROM : GOMP_MAP_TOFROM; else if (strcmp ("release", p) == 0) kind = GOMP_MAP_RELEASE; else if (strcmp ("delete", p) == 0) @@ -15735,35 +15764,6 @@ c_parser_omp_clause_map (c_parser *parser, tree list) c_parser_consume_token (parser); c_parser_consume_token (parser); } - else if (always) - { - if (always_id_kind != C_ID_ID) - { - c_parser_error (parser, "expected identifier"); - parens.skip_until_found_close (parser); - return list; - } - - tree t = lookup_name (always_id); - if (t == NULL_TREE) - { - undeclared_variable (always_loc, always_id); - t = error_mark_node; - } - if (t != error_mark_node) - { - tree u = build_omp_clause (clause_loc, OMP_CLAUSE_MAP); - OMP_CLAUSE_DECL (u) = t; - OMP_CLAUSE_CHAIN (u) = list; - OMP_CLAUSE_SET_MAP_KIND (u, kind); - list = u; - } - if (always == 1) - { - parens.skip_until_found_close (parser); - return list; - } - } nl = c_parser_omp_variable_list (parser, clause_loc, OMP_CLAUSE_MAP, list, C_ORT_OMP, true); diff --git a/gcc/cp/ChangeLog.omp b/gcc/cp/ChangeLog.omp index 8e6fa988467..3718a548610 100644 --- a/gcc/cp/ChangeLog.omp +++ b/gcc/cp/ChangeLog.omp @@ -1,3 +1,11 @@ +2021-05-14 Tobias Burnus + + Backported from master: + 2021-05-12 Marcel Vollweiler + + * parser.c (cp_parser_omp_clause_map): Support map-type-modifier + 'close'. + 2021-05-14 Tobias Burnus Backported from master: diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index 8910ebe5124..6da9cdb9392 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -37859,40 +37859,90 @@ cp_parser_omp_clause_depend (cp_parser *parser, tree list, location_t loc) map-kind: alloc | to | from | tofrom | release | delete - map ( always [,] map-kind: variable-list ) */ + map ( always [,] map-kind: variable-list ) + + OpenMP 5.0: + map ( [map-type-modifier[,] ...] map-kind: variable-list ) + + map-type-modifier: + always | close */ static tree cp_parser_omp_clause_map (cp_parser *parser, tree list) { tree nlist, c; enum gomp_map_kind kind = GOMP_MAP_TOFROM; - bool always = false; if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN)) return list; - if (cp_lexer_next_token_is (parser->lexer, CPP_NAME)) + int pos = 1; + int map_kind_pos = 0; + while (cp_lexer_peek_nth_token (parser->lexer, pos)->type == CPP_NAME + || cp_lexer_peek_nth_token (parser->lexer, pos)->keyword == RID_DELETE) { - tree id = cp_lexer_peek_token (parser->lexer)->u.value; - const char *p = IDENTIFIER_POINTER (id); + if (cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type == CPP_COLON) + { + map_kind_pos = pos; + break; + } + + if (cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type == CPP_COMMA) + pos++; + pos++; + } + bool always_modifier = false; + bool close_modifier = false; + for (int pos = 1; pos < map_kind_pos; ++pos) + { + cp_token *tok = cp_lexer_peek_token (parser->lexer); + if (tok->type == CPP_COMMA) + { + cp_lexer_consume_token (parser->lexer); + continue; + } + + const char *p = IDENTIFIER_POINTER (tok->u.value); if (strcmp ("always", p) == 0) { - int nth = 2; - if (cp_lexer_peek_nth_token (parser->lexer, 2)->type == CPP_COMMA) - nth++; - if ((cp_lexer_peek_nth_token (parser->lexer, nth)->type == CPP_NAME - || (cp_lexer_peek_nth_token (parser->lexer, nth)->keyword - == RID_DELETE)) - && (cp_lexer_peek_nth_token (parser->lexer, nth + 1)->type - == CPP_COLON)) + if (always_modifier) { - always = true; - cp_lexer_consume_token (parser->lexer); - if (nth == 3) - cp_lexer_consume_token (parser->lexer); + cp_parser_error (parser, "too many % modifiers"); + cp_parser_skip_to_closing_parenthesis (parser, + /*recovering=*/true, + /*or_comma=*/false, + /*consume_paren=*/true); + return list; + } + always_modifier = true; + } + else if (strcmp ("close", p) == 0) + { + if (close_modifier) + { + cp_parser_error (parser, "too many % modifiers"); + cp_parser_skip_to_closing_parenthesis (parser, + /*recovering=*/true, + /*or_comma=*/false, + /*consume_paren=*/true); + return list; } + close_modifier = true; } + else + { + cp_parser_error (parser, "%<#pragma omp target%> with " + "modifier other than % or %" + "on % clause"); + cp_parser_skip_to_closing_parenthesis (parser, + /*recovering=*/true, + /*or_comma=*/false, + /*consume_paren=*/true); + return list; + } + + cp_lexer_consume_token (parser->lexer); } if (cp_lexer_next_token_is (parser->lexer, CPP_NAME) @@ -37904,11 +37954,11 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list) if (strcmp ("alloc", p) == 0) kind = GOMP_MAP_ALLOC; else if (strcmp ("to", p) == 0) - kind = always ? GOMP_MAP_ALWAYS_TO : GOMP_MAP_TO; + kind = always_modifier ? GOMP_MAP_ALWAYS_TO : GOMP_MAP_TO; else if (strcmp ("from", p) == 0) - kind = always ? GOMP_MAP_ALWAYS_FROM : GOMP_MAP_FROM; + kind = always_modifier ? GOMP_MAP_ALWAYS_FROM : GOMP_MAP_FROM; else if (strcmp ("tofrom", p) == 0) - kind = always ? GOMP_MAP_ALWAYS_TOFROM : GOMP_MAP_TOFROM; + kind = always_modifier ? GOMP_MAP_ALWAYS_TOFROM : GOMP_MAP_TOFROM; else if (strcmp ("release", p) == 0) kind = GOMP_MAP_RELEASE; else diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp index 01b68bdf790..040716a0096 100644 --- a/gcc/testsuite/ChangeLog.omp +++ b/gcc/testsuite/ChangeLog.omp @@ -1,3 +1,11 @@ +2021-05-14 Tobias Burnus + + Backported from master: + 2021-05-12 Marcel Vollweiler + + * c-c++-common/gomp/map-6.c: New test. + * c-c++-common/gomp/map-7.c: New test. + 2021-05-14 Tobias Burnus Backported from master: diff --git a/gcc/testsuite/c-c++-common/gomp/map-6.c b/gcc/testsuite/c-c++-common/gomp/map-6.c new file mode 100644 index 00000000000..6ee59714847 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/map-6.c @@ -0,0 +1,135 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-original" } */ + +void +foo (void) +{ + /* Test to ensure that the close modifier is parsed and ignored in map clauses. */ + int a, b, b1, b2, b3, b4, b5, b6, b7; + + #pragma omp target map (a) + ; + + #pragma omp target map (to:a) + ; + + #pragma omp target map (a to: b) /* { dg-error "'#pragma omp target' with modifier other than 'always' or 'close'" } */ + ; + + #pragma omp target map (close, a to: b) /* { dg-error "'#pragma omp target' with modifier other than 'always' or 'close'" } */ + ; + + #pragma omp target map (close a) /* { dg-error "'close' undeclared" "" { target c } } */ + /* { dg-error "'close' has not been declared" "" { target c++ } .-1 } */ + /* { dg-error "expected '\\)' before 'a'" "" { target *-*-* } .-2 } */ + ; + + #pragma omp target map (always a) /* { dg-error "'always' undeclared" "" { target c } } */ + /* { dg-error "'always' has not been declared" "" { target c++ } .-1 } */ + /* { dg-error "expected '\\)' before 'a'" "" { target *-*-* } .-2 } */ + ; + + #pragma omp target map (close to:a) + ; + + #pragma omp target map (close, to:a) + ; + + #pragma omp target map (close delete:a) /* { dg-error "'#pragma omp target' with map-type other than 'to', 'from', 'tofrom' or 'alloc' on 'map' clause" } */ + ; + + #pragma omp target map (close always to:b1) + ; + + #pragma omp target map (close, always to:b2) + ; + + #pragma omp target map (close, always, to:b3) + ; + + #pragma omp target map (always close to:b4) + ; + + #pragma omp target map (always, close to:b5) + ; + + #pragma omp target map (always, close, to:b6) + ; + + #pragma omp target map (always, always, to:a) /* { dg-error "too many 'always' modifiers" } */ + ; + + #pragma omp target map (always always, to:a) /* { dg-error "too many 'always' modifiers" } */ + ; + + #pragma omp target map (always, always to:a) /* { dg-error "too many 'always' modifiers" } */ + ; + + #pragma omp target map (always always to:a) /* { dg-error "too many 'always' modifiers" } */ + ; + + #pragma omp target map (close, close, to:a) /* { dg-error "too many 'close' modifiers" } */ + ; + + #pragma omp target map (close close, to:a) /* { dg-error "too many 'close' modifiers" } */ + ; + + #pragma omp target map (close, close to:a) /* { dg-error "too many 'close' modifiers" } */ + ; + + #pragma omp target map (close close to:a) /* { dg-error "too many 'close' modifiers" } */ + ; + + #pragma omp target map (always to : a) map (close to : b) + ; + + int close = 0; + #pragma omp target map (close) + ; + + #pragma omp target map (close a) /* { dg-error "expected '\\)' before 'a'" } */ + ; + + int always = 0; + #pragma omp target map (always) + ; + + #pragma omp target map (always a) /* { dg-error "expected '\\)' before 'a'" } */ + ; + + #pragma omp target map (always, close) + ; + + #pragma omp target map (always, always) /* { dg-error "'always' appears more than once in map clauses" } */ + ; + + #pragma omp target map (always, always, close) /* { dg-error "'always' appears more than once in map clauses" } */ + ; + + #pragma omp target map (always, close, to: always, close, b7) + ; + + int to = 0; + #pragma omp target map (always, close, to) + ; + + #pragma omp target map (to, always, close) + { + to = always = close = 1; + } + if (to != 1 || always != 1 || close != 1) + __builtin_abort (); + ; +} + +/* { dg-final { scan-tree-dump-not "map\\(\[^\n\r)]*close\[^\n\r)]*to:" "original" } } */ + +/* { dg-final { scan-tree-dump-times "pragma omp target map\\(always,to:" 7 "original" } } */ + +/* { dg-final { scan-tree-dump "pragma omp target map\\(always,to:b1" "original" } } */ +/* { dg-final { scan-tree-dump "pragma omp target map\\(always,to:b2" "original" } } */ +/* { dg-final { scan-tree-dump "pragma omp target map\\(always,to:b3" "original" } } */ +/* { dg-final { scan-tree-dump "pragma omp target map\\(always,to:b4" "original" } } */ +/* { dg-final { scan-tree-dump "pragma omp target map\\(always,to:b5" "original" } } */ +/* { dg-final { scan-tree-dump "pragma omp target map\\(always,to:b6" "original" } } */ +/* { dg-final { scan-tree-dump "pragma omp target map\\(always,to:b7\\) map\\(always,to:close\\) map\\(always,to:always\\)" "original" } } */ diff --git a/gcc/testsuite/c-c++-common/gomp/map-7.c b/gcc/testsuite/c-c++-common/gomp/map-7.c new file mode 100644 index 00000000000..3f1e9729e41 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/map-7.c @@ -0,0 +1,20 @@ +/* { dg-do compile } */ + +void +foo (void) +{ + /* Test to ensure that the close modifier is parsed and ignored in map clauses. */ + + #define N 1024 + int always[N]; + int close; + + #pragma omp target map(always[:N]) + ; + + #pragma omp target map(close, always[:N]) + ; + + #pragma omp target map(always[:N], close) + ; +}