From patchwork Mon Jul 10 18:33:58 2023 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Chung-Lin Tang X-Patchwork-Id: 1805969 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@legolas.ozlabs.org Authentication-Results: legolas.ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=2620:52:3:1:0:246e:9693:128c; helo=server2.sourceware.org; envelope-from=gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Authentication-Results: legolas.ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.a=rsa-sha256 header.s=default header.b=lShR7DPM; dkim-atps=neutral Received: from server2.sourceware.org (server2.sourceware.org [IPv6:2620:52:3:1:0:246e:9693:128c]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature ECDSA (P-384) server-digest SHA384) (No client certificate requested) by legolas.ozlabs.org (Postfix) with ESMTPS id 4R0CMr61vtz20Ph for ; Tue, 11 Jul 2023 04:34:31 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id C7E953857704 for ; Mon, 10 Jul 2023 18:34:29 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org C7E953857704 DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gcc.gnu.org; s=default; t=1689014069; bh=ZFkDQgsRVM+L8RrBvjZ7hmlDea0GHNJo9NBwwc6ojqo=; h=Date:Subject:To:List-Id:List-Unsubscribe:List-Archive:List-Post: List-Help:List-Subscribe:From:Reply-To:Cc:From; b=lShR7DPMcYxqbJuQDgZf2Z431Md5i17KuDgwPkrOJ2ktkdQ3UiEcRJGDxrnaGr3GK V5/Cu/fR8X51ubzVoodCEm14Wr9nzwl6Z995zfNpN+4IHqse2RDykMPMmHqo94kMls npeZWfxmJMJfD3709nlJ07ne/kawBUhqdg/IIPLw= X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from APC01-PSA-obe.outbound.protection.outlook.com (mail-psaapc01on2070.outbound.protection.outlook.com [40.107.255.70]) by sourceware.org (Postfix) with ESMTPS id 6196F3858D32 for ; Mon, 10 Jul 2023 18:34:08 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 6196F3858D32 ARC-Seal: i=1; a=rsa-sha256; s=arcselector9901; d=microsoft.com; cv=none; b=ZbYY4QLa//xJnQHJAzP+CnDh2RZLKgb1keQztIYcqhqLq2oT2fK7AVHaitLXKKEhWZ1WNuTBVZ7cKEWdbLBpqNa5FGZChqiwL/mLAVwW7kFWi8DdX3yUfJoi3RGFYknA6km+WSJWbQew0D7a1RCH1Y9tz3IBqiw01AAuhF8GeYBIFGKF8k6FP2uzq1eRKzSB3BYvigz/8ftVzj6xdACbSbydYQA3uOGDcWRMUKj7VwBlbR2r41JX2RN9sLVNyHQrfuLBDJCcj1HgGZpAR3lvuNm1fc1gjoFxNQPgAsLCNPepCQtFWf0MvpaMUDVB6y63lruxm7Z8ObdvT1Bn/HqzRA== ARC-Message-Signature: i=1; a=rsa-sha256; c=relaxed/relaxed; d=microsoft.com; s=arcselector9901; h=From:Date:Subject:Message-ID:Content-Type:MIME-Version:X-MS-Exchange-AntiSpam-MessageData-ChunkCount:X-MS-Exchange-AntiSpam-MessageData-0:X-MS-Exchange-AntiSpam-MessageData-1; bh=ZFkDQgsRVM+L8RrBvjZ7hmlDea0GHNJo9NBwwc6ojqo=; b=dhRlGVYTdXa74pqyzTBI5l2J/vu6G5HtXOnuE6wyfyqp6sG1+vEcDdZaLaiYtXxoUNfjFFgxzNTwbfZbV1CYYRNE37VeHNUrywQKUgPAA/aBLmeICgqy4BhNTWfowifGkoz5Btib7QPA01W7j/D3WscuLGUP8Qt0EdGZaEb8ViDGeP3StBT+xGiCHhXtWNgRTItmFnbot3yQNuf8v/vQvp6+BoD7sEXH8iTUD92WelAY4mV1tE3N4cq2iXzLDjWhHaf318dEyhLEm9Lv4ntp+3TLdEdcuJ0LHj//4cStO5iI3KLqQoxXjwgNHgeALqwXshVZRkqJVbNI/it4/LCOZQ== ARC-Authentication-Results: i=1; mx.microsoft.com 1; spf=pass smtp.mailfrom=siemens.com; dmarc=pass action=none header.from=siemens.com; dkim=pass header.d=siemens.com; arc=none Received: from SG2PR06MB5430.apcprd06.prod.outlook.com (2603:1096:4:1ba::14) by SEZPR06MB5141.apcprd06.prod.outlook.com (2603:1096:101:43::14) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_256_GCM_SHA384) id 15.20.6565.30; Mon, 10 Jul 2023 18:34:03 +0000 Received: from SG2PR06MB5430.apcprd06.prod.outlook.com ([fe80::e9fc:a1ea:cfed:90b8]) by SG2PR06MB5430.apcprd06.prod.outlook.com ([fe80::e9fc:a1ea:cfed:90b8%6]) with mapi id 15.20.6565.028; Mon, 10 Jul 2023 18:34:02 +0000 Message-ID: Date: Tue, 11 Jul 2023 02:33:58 +0800 User-Agent: Mozilla/5.0 (Macintosh; Intel Mac OS X 10.15; rv:102.0) Gecko/20100101 Thunderbird/102.13.0 Content-Language: en-US Subject: [PATCH, OpenACC 2.7] readonly modifier support in front-ends To: gcc-patches , Thomas Schwinge , Catherine Moore , Tobias Burnus X-ClientProxiedBy: TY2PR0101CA0019.apcprd01.prod.exchangelabs.com (2603:1096:404:92::31) To SG2PR06MB5430.apcprd06.prod.outlook.com (2603:1096:4:1ba::14) MIME-Version: 1.0 X-MS-PublicTrafficType: Email X-MS-TrafficTypeDiagnostic: SG2PR06MB5430:EE_|SEZPR06MB5141:EE_ X-MS-Office365-Filtering-Correlation-Id: fe46458b-7807-4008-c130-08db81743b78 X-MS-Exchange-SenderADCheck: 1 X-MS-Exchange-AntiSpam-Relay: 0 X-Microsoft-Antispam: BCL:0; X-Microsoft-Antispam-Message-Info: ZuE2+AXd73bu0PG3mQ6euXr1lleUCsvjo7DsauNa/xf5E5vPCuP9d4v7wlvLA93/ZwH8v2z+Trz83rALxulwONcRszY7DH3YeqEguLJ9puXJH8TtNkysRTsAQAAKRkIhR34qUNvnz3VCwnHdvjQSAQ/FYj6NuGjt/30p40XfJ7/4B1SsqRpahBzjcWKG50bt/E+HyoeEDS3xHQVgh4ctPpkI6W+IKss9sw16wBvLYlUorWAeWmvQ/Q2f2Bme8OPhutWnW/57BCqsDsOW4Mw+sBBJXKOiRL+EoXAiFDiOsNv60dFs/BnAlqwmDT6xZbSRU1FtLIGnQrPnbmoeMEI56ok6v1cV5+yA52IJ+Z3+QSTHR3c/6dQlEe310t2Va8vAUfum1HaGNzG1rnNo4gkl9BtFf4/kjEB/nsAidghNKtOqy+ErPgHwe29OjKc0hotno3iUjJPUP1e6z3vCDoBxtQn6p9RNAUwtgWerMZ33d9fkhsYgfN0XcIcM3vOLZWwxijhdyRvgtNk9Zv4mRVvYGHOmObQdbefYmj74dpdCn8qAqX4uwylDhJwO6kevl6W3IFsS/ZYSnqyLpzfU5LKGLbinRZqP6j7OVhTqD3jtkiUXM8gF6SukU81Xy8mVICad7UQNPjW+3o8dYpraNpAoJQ== X-Forefront-Antispam-Report: CIP:255.255.255.255; CTRY:; LANG:en; SCL:1; SRV:; IPV:NLI; SFV:NSPM; H:SG2PR06MB5430.apcprd06.prod.outlook.com; PTR:; CAT:NONE; SFS:(13230028)(4636009)(39860400002)(136003)(346002)(366004)(396003)(376002)(451199021)(186003)(26005)(2616005)(6512007)(6506007)(83380400001)(41300700001)(66556008)(316002)(2906002)(235185007)(5660300002)(66476007)(8676002)(66946007)(8936002)(478600001)(6486002)(33964004)(6666004)(110136005)(36756003)(82960400001)(31696002)(38100700002)(86362001)(84970400001)(31686004)(43740500002)(45980500001); DIR:OUT; SFP:1101; X-MS-Exchange-AntiSpam-MessageData-ChunkCount: 1 X-MS-Exchange-AntiSpam-MessageData-0: =?utf-8?q?Zr/pJwT8ZN0vm3DuP1iYXfoOA1gT?= =?utf-8?q?dMDsPuwikDjuQdocsgmpgS20Qy+nYAQYte+GG26kR/Gq4gVGNOXZUuSX80VVM4F1i?= =?utf-8?q?eP/vG0AiVEknjEGhwihd207Z+P8cGWUv4CpA5WlPfonDF12UUD4KPNMwiD61/BQuq?= =?utf-8?q?nRl3XbbCxijt0PbUEKvS3vx9TIrB49wWb3ivCODI5saCWTYtCQp2I/FrQSNvh1BG9?= =?utf-8?q?fi72F618qQwIwROtRrYWaT4set91RXrkQdtQZTnKH55/h5DKXbGol58zbHcxN8wIa?= =?utf-8?q?YMeameBmaQcuJXvcyJ9A8PclYaUm7Ie8K/RWsJe0uiRY9dhu3QxtT/C5PBl6dK1jO?= =?utf-8?q?UviXHMwlLAeKB2jhVFUGR3csDlvOaPncxOB6vusI18ZXax3n+jNffIr3W8Bit8G16?= =?utf-8?q?+jk7VqpWOZ1XIWnJtcZyeJVgz9s+BrEy+mxFj+QrHPtKNORCZrpzWUFip7PY9xfYI?= =?utf-8?q?/5MTi1SrpRDPco50g/Jh6pITVLU/7x+bXWKukAVIe0tgp4L0LPfHEBldR6sN4MX2n?= =?utf-8?q?PJfc5L9IRIoL94R43dIEikSLxzVUJRKQjhO+oPq/2OnktMfykWh+jSLylvv0MsZeF?= =?utf-8?q?fq3aeBBv3+yLLZzYfbZyv/wyQtC2Ydrr+uUTykl4+aNQc2QP/PAYhvnrdsm/alZkV?= =?utf-8?q?jr31r80Mg3r9F4VIEmzUzkYH+/iaIt6x2bS0Cb1bTH0FP+yBn2KGU5e3kgDDDpMhM?= =?utf-8?q?sCavlCuZqLidvHicEU9eFU1sEixIdyE8QkhcvIceAd4MvpWYkDxo5zUZvqyZBvNJ9?= =?utf-8?q?lU2ongR6XS0JDJogsq5eUfVCpTZZefxwcAZ4kQ5lM5FKlhDanj205jN7321u6qpfM?= =?utf-8?q?zgtre7jLCtbJbOTxYr4cqWdeLYPWVjT3yP7pza1SfRqsi8bBxlmcn3pSHlzfr/Y0y?= =?utf-8?q?JtMKVzyo5Ecjamb3ReU0AWRaSb9drH1uD7yDkErP5I9jSV2Hc84aogavaaDt5Prqh?= =?utf-8?q?irJprmEr65Iyh7cI8frAiG8H8FpNjJSM1pe7deXtzHehJ9OrAXSLHRb1dqU6x6EgU?= =?utf-8?q?XPbLFy7A4yKkYA1PyBddGUK5EB8BODoS8yzN5k7Nr1Cv7HDj4vcD9TCxhjy2cN6dl?= =?utf-8?q?akyfvD/4lq/quZWeSjF20CTpPXLFsXl06V9cyT98gZ2nmSfIKUuL3bsb3kN94lzWy?= =?utf-8?q?QnSR2tpnweUhwFOvZ6UdvPBziy5W2fk2NROBsAD6DLqEpq49XKUwFoK+e0KNqq9oZ?= =?utf-8?q?FG/112Z4SDy5uDRJBwt4bEtEQY3KTMEgjGRypE/WDH+4MquECRrwZvsWgai26T18z?= =?utf-8?q?g+d1zmabFPjweXeUu8z0NlNEN5+UnHKanAbjGcYp6ZmFqo1pNJwSvVNC61OZiZV79?= =?utf-8?q?WTPsXPHvONh6rskTkGNLdqva/YMRfmL9XH4GuPxUh3wQ/Fahzkvj1Du0SkFp9iTfO?= =?utf-8?q?8UnH56NSyD2fgDsOiGdWKrhFMKDfc3rAIBYOo14ah9Ha+0epGFmCzppZHfo6EGmNr?= =?utf-8?q?gFeH/2XH5gpXDLc/3OAN44voNztB8XGhWbBMOuJiiFBbkjU/iuJmAesY6pX5rrveP?= =?utf-8?q?R/lUhBd6FKKEbUG46ILthceZIVpGk6+gMg=3D=3D?= X-OriginatorOrg: siemens.com X-MS-Exchange-CrossTenant-Network-Message-Id: fe46458b-7807-4008-c130-08db81743b78 X-MS-Exchange-CrossTenant-AuthSource: SG2PR06MB5430.apcprd06.prod.outlook.com X-MS-Exchange-CrossTenant-AuthAs: Internal X-MS-Exchange-CrossTenant-OriginalArrivalTime: 10 Jul 2023 18:34:02.7321 (UTC) X-MS-Exchange-CrossTenant-FromEntityHeader: Hosted X-MS-Exchange-CrossTenant-Id: 38ae3bcd-9579-4fd4-adda-b42e1495d55a X-MS-Exchange-CrossTenant-MailboxType: HOSTED X-MS-Exchange-CrossTenant-UserPrincipalName: 097K0PLdvH6oZxEtbQ0n0+jxcF6rkKnIWWz8bMlj2aOxAjZg+DPlYHzaG4iA4w8dJogjjHcD12Ax+4/P5agv7O6uGz5fuGSVQNvcHJt5AsA= X-MS-Exchange-Transport-CrossTenantHeadersStamped: SEZPR06MB5141 X-Spam-Status: No, score=-9.7 required=5.0 tests=BAYES_00, DKIMWL_WL_MED, DKIM_SIGNED, DKIM_VALID, DKIM_VALID_AU, DKIM_VALID_EF, FORGED_SPF_HELO, GIT_PATCH_0, RCVD_IN_DNSWL_NONE, RCVD_IN_MSPIKE_H2, SPF_HELO_PASS, SPF_NONE, TXREP, T_SCC_BODY_TEXT_LINE autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , X-Patchwork-Original-From: Chung-Lin Tang via Gcc-patches From: Chung-Lin Tang Reply-To: cltang@codesourcery.com Cc: Chung-Lin Tang Errors-To: gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org Sender: "Gcc-patches" Hi Thomas, this patch contains support for the 'readonly' modifier in copyin clauses and the cache directive. As we discussed earlier, the work for actually linking this to middle-end points-to analysis is a somewhat non-trivial issue. This first patch allows the language feature to be used in OpenACC directives first (with no effect for now). The middle-end changes are probably going to be a later patch. (Also CCing Tobias because of the Fortran bits) Tested on powerpc64le-linux with nvptx offloading. Is this okay for trunk? Thanks, Chung-Lin 2023-07-10 Chung-Lin Tang gcc/c/ChangeLog: * c-parser.cc (c_parser_omp_var_list_parens): Add 'bool *readonly = NULL' parameter, add readonly modifier parsing support. (c_parser_oacc_data_clause): Adjust c_parser_omp_var_list_parens call to turn on readonly modifier parsing for copyin clause, set OMP_CLAUSE_MAP_READONLY if readonly modifier found, update comments. (c_parser_oacc_cache): Adjust c_parser_omp_var_list_parens call to turn on readonly modifier parsing, set OMP_CLAUSE__CACHE__READONLY if readonly modifier found, update comments. gcc/cp/ChangeLog: * parser.cc (cp_parser_omp_var_list): Add 'bool *readonly = NULL' parameter, add readonly modifier parsing support. (cp_parser_oacc_data_clause): Adjust cp_parser_omp_var_list call to turn on readonly modifier parsing for copyin clause, set OMP_CLAUSE_MAP_READONLY if readonly modifier found, update comments. (cp_parser_oacc_cache): Adjust cp_parser_omp_var_list call to turn on readonly modifier parsing, set OMP_CLAUSE__CACHE__READONLY if readonly modifier found, update comments. gcc/fortran/ChangeLog: * gfortran.h (typedef struct gfc_omp_namelist): Adjust map_op as ENUM_BITFIELD field, add 'bool readonly' field. * openmp.cc (gfc_match_omp_map_clause): Add 'bool readonly = false' parameter, set n->u.readonly field. (gfc_match_omp_clauses): Add readonly modifier parsing for OpenACC copyin clause, adjust call to gfc_match_omp_map_clause. (gfc_match_oacc_cache): Add readonly modifier parsing for OpenACC cache directive, adjust call to gfc_match_omp_map_clause. * trans-openmp.cc (gfc_trans_omp_clauses): Set OMP_CLAUSE_MAP_READONLY, OMP_CLAUSE__CACHE__READONLY to 1 when readonly is set. gcc/ChangeLog: * tree-pretty-print.cc (dump_omp_clause): Add support for printing OMP_CLAUSE_MAP_READONLY and OMP_CLAUSE__CACHE__READONLY. * tree.h (OMP_CLAUSE_MAP_READONLY): New macro. (OMP_CLAUSE__CACHE__READONLY): New macro. gcc/testsuite/ChangeLog: * c-c++-common/goacc/readonly-1.c: New test. * gfortran.dg/goacc/readonly-1.f90: New test. diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc index d4b98d5d8b6..09e1e89d793 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -14059,7 +14059,8 @@ c_parser_omp_variable_list (c_parser *parser, static tree c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind, - tree list, bool allow_deref = false) + tree list, bool allow_deref = false, + bool *readonly = NULL) { /* The clauses location. */ location_t loc = c_parser_peek_token (parser)->location; @@ -14067,6 +14068,20 @@ c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind, matching_parens parens; if (parens.require_open (parser)) { + if (readonly != NULL) + { + c_token *token = c_parser_peek_token (parser); + if (token->type == CPP_NAME + && !strcmp (IDENTIFIER_POINTER (token->value), "readonly") + && c_parser_peek_2nd_token (parser)->type == CPP_COLON) + { + c_parser_consume_token (parser); + c_parser_consume_token (parser); + *readonly = true; + } + else + *readonly = false; + } list = c_parser_omp_variable_list (parser, loc, kind, list, allow_deref); parens.skip_until_found_close (parser); } @@ -14084,7 +14099,11 @@ c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind, OpenACC 2.6: no_create ( variable-list ) attach ( variable-list ) - detach ( variable-list ) */ + detach ( variable-list ) + + OpenACC 2.7: + copyin (readonly : variable-list ) + */ static tree c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind, @@ -14135,11 +14154,22 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind, default: gcc_unreachable (); } + + /* Turn on readonly modifier parsing for copyin clause. */ + bool readonly = false, *readonly_ptr = NULL; + if (c_kind == PRAGMA_OACC_CLAUSE_COPYIN) + readonly_ptr = &readonly; + tree nl, c; - nl = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_MAP, list, true); + nl = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_MAP, list, true, + readonly_ptr); for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c)) - OMP_CLAUSE_SET_MAP_KIND (c, kind); + { + OMP_CLAUSE_SET_MAP_KIND (c, kind); + if (readonly) + OMP_CLAUSE_MAP_READONLY (c) = 1; + } return nl; } @@ -18212,6 +18242,9 @@ c_parser_omp_structured_block (c_parser *parser, bool *if_p) /* OpenACC 2.0: # pragma acc cache (variable-list) new-line + OpenACC 2.7: + # pragma acc cache (readonly: variable-list) new-line + LOC is the location of the #pragma token. */ @@ -18219,8 +18252,14 @@ static tree c_parser_oacc_cache (location_t loc, c_parser *parser) { tree stmt, clauses; + bool readonly; + + clauses = c_parser_omp_var_list_parens (parser, OMP_CLAUSE__CACHE_, NULL, + false, &readonly); + if (readonly) + for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) + OMP_CLAUSE__CACHE__READONLY (c) = 1; - clauses = c_parser_omp_var_list_parens (parser, OMP_CLAUSE__CACHE_, NULL); clauses = c_finish_omp_clauses (clauses, C_ORT_ACC); c_parser_skip_to_pragma_eol (parser); diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc index acd1bd48af5..0f51289539b 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -37727,11 +37727,27 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind, static tree cp_parser_omp_var_list (cp_parser *parser, enum omp_clause_code kind, tree list, - bool allow_deref = false) + bool allow_deref = false, bool *readonly = NULL) { if (cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN)) - return cp_parser_omp_var_list_no_open (parser, kind, list, NULL, - allow_deref); + { + if (readonly != NULL) + { + cp_token *token = cp_lexer_peek_token (parser->lexer); + if (token->type == CPP_NAME + && !strcmp (IDENTIFIER_POINTER (token->u.value), "readonly") + && cp_lexer_nth_token_is (parser->lexer, 2, CPP_COLON)) + { + cp_lexer_consume_token (parser->lexer); + cp_lexer_consume_token (parser->lexer); + *readonly = true; + } + else + *readonly = false; + } + return cp_parser_omp_var_list_no_open (parser, kind, list, NULL, + allow_deref); + } return list; } @@ -37746,7 +37762,11 @@ cp_parser_omp_var_list (cp_parser *parser, enum omp_clause_code kind, tree list, OpenACC 2.6: no_create ( variable-list ) attach ( variable-list ) - detach ( variable-list ) */ + detach ( variable-list ) + + OpenACC 2.7: + copyin (readonly : variable-list ) + */ static tree cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind, @@ -37797,11 +37817,22 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind, default: gcc_unreachable (); } + + /* Turn on readonly modifier parsing for copyin clause. */ + bool readonly = false, *readonly_ptr = NULL; + if (c_kind == PRAGMA_OACC_CLAUSE_COPYIN) + readonly_ptr = &readonly; + tree nl, c; - nl = cp_parser_omp_var_list (parser, OMP_CLAUSE_MAP, list, true); + nl = cp_parser_omp_var_list (parser, OMP_CLAUSE_MAP, list, true, + readonly_ptr); for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c)) - OMP_CLAUSE_SET_MAP_KIND (c, kind); + { + OMP_CLAUSE_SET_MAP_KIND (c, kind); + if (readonly) + OMP_CLAUSE_MAP_READONLY (c) = 1; + } return nl; } @@ -45875,6 +45906,9 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok, /* OpenACC 2.0: # pragma acc cache (variable-list) new-line + + OpenACC 2.7: + # pragma acc cache (readonly: variable-list) new-line */ static tree @@ -45885,8 +45919,14 @@ cp_parser_oacc_cache (cp_parser *parser, cp_token *pragma_tok) auto_suppress_location_wrappers sentinel; tree stmt, clauses; + bool readonly; + + clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE__CACHE_, NULL, + false, &readonly); + if (readonly) + for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) + OMP_CLAUSE__CACHE__READONLY (c) = 1; - clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE__CACHE_, NULL_TREE); clauses = finish_omp_clauses (clauses, C_ORT_ACC); cp_parser_require_pragma_eol (parser, cp_lexer_peek_token (parser->lexer)); diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h index cc7ba7c8846..9fa8962d63f 100644 --- a/gcc/fortran/gfortran.h +++ b/gcc/fortran/gfortran.h @@ -1360,7 +1360,11 @@ typedef struct gfc_omp_namelist { gfc_omp_reduction_op reduction_op; gfc_omp_depend_doacross_op depend_doacross_op; - gfc_omp_map_op map_op; + struct + { + ENUM_BITFIELD (gfc_omp_map_op) map_op:8; + bool readonly; + }; gfc_expr *align; struct { diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc index 038907baa48..acd1428d2d7 100644 --- a/gcc/fortran/openmp.cc +++ b/gcc/fortran/openmp.cc @@ -1196,7 +1196,7 @@ omp_inv_mask::omp_inv_mask (const omp_mask &m) : omp_mask (m) static bool gfc_match_omp_map_clause (gfc_omp_namelist **list, gfc_omp_map_op map_op, - bool allow_common, bool allow_derived) + bool allow_common, bool allow_derived, bool readonly = false) { gfc_omp_namelist **head = NULL; if (gfc_match_omp_variable_list ("", list, allow_common, NULL, &head, true, @@ -1205,7 +1205,10 @@ gfc_match_omp_map_clause (gfc_omp_namelist **list, gfc_omp_map_op map_op, { gfc_omp_namelist *n; for (n = *head; n; n = n->next) - n->u.map_op = map_op; + { + n->u.map_op = map_op; + n->u.readonly = readonly; + } return true; } @@ -2079,11 +2082,16 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, { if (openacc) { - if (gfc_match ("copyin ( ") == MATCH_YES - && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_TO, true, - allow_derived)) - continue; + if (gfc_match ("copyin ( ") == MATCH_YES) + { + bool readonly = false; + if (gfc_match ("readonly : ") == MATCH_YES) + readonly = true; + if (gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], + OMP_MAP_TO, true, + allow_derived, readonly)) + continue; + } } else if (gfc_match_omp_variable_list ("copyin (", &c->lists[OMP_LIST_COPYIN], @@ -4008,20 +4016,35 @@ gfc_match_oacc_wait (void) match gfc_match_oacc_cache (void) { + bool readonly = false; gfc_omp_clauses *c = gfc_get_omp_clauses (); /* The OpenACC cache directive explicitly only allows "array elements or subarrays", which we're currently not checking here. Either check this after the call of gfc_match_omp_variable_list, or add something like a only_sections variant next to its allow_sections parameter. */ - match m = gfc_match_omp_variable_list (" (", - &c->lists[OMP_LIST_CACHE], true, - NULL, NULL, true); + match m = gfc_match (" ( "); if (m != MATCH_YES) { gfc_free_omp_clauses(c); return m; } + if (gfc_match ("readonly :") == MATCH_YES) + readonly = true; + + gfc_omp_namelist **head = NULL; + m = gfc_match_omp_variable_list ("", &c->lists[OMP_LIST_CACHE], true, + NULL, &head, true); + if (m != MATCH_YES) + { + gfc_free_omp_clauses(c); + return m; + } + + if (readonly) + for (gfc_omp_namelist *n = *head; n; n = n->next) + n->u.readonly = true; + if (gfc_current_state() != COMP_DO && gfc_current_state() != COMP_DO_CONCURRENT) { diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc index 0f8323901d7..87d0b5e0cdf 100644 --- a/gcc/fortran/trans-openmp.cc +++ b/gcc/fortran/trans-openmp.cc @@ -3067,6 +3067,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, || (n->expr && gfc_expr_attr (n->expr).pointer))) always_modifier = true; + if (n->u.readonly) + OMP_CLAUSE_MAP_READONLY (node) = 1; + switch (n->u.map_op) { case OMP_MAP_ALLOC: @@ -3920,6 +3923,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, } if (n->u.present_modifier) OMP_CLAUSE_MOTION_PRESENT (node) = 1; + if (list == OMP_LIST_CACHE && n->u.readonly) + OMP_CLAUSE__CACHE__READONLY (node) = 1; omp_clauses = gfc_trans_add_clause (node, omp_clauses); } break; diff --git a/gcc/testsuite/c-c++-common/goacc/readonly-1.c b/gcc/testsuite/c-c++-common/goacc/readonly-1.c new file mode 100644 index 00000000000..171f96c08db --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/readonly-1.c @@ -0,0 +1,27 @@ +/* { dg-additional-options "-fdump-tree-original" } */ + +struct S +{ + int *ptr; + float f; +}; + + +int main (void) +{ + int x[32]; + struct S s = {x, 0}; + + #pragma acc parallel copyin(readonly: x[:32], s.ptr[:16]) + { + #pragma acc cache (readonly: x[:32]) + } + return 0; +} + +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:\\*s.ptr \\\[len: 64\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: 128\\\]\\)" 1 "original" { target { c } } } } */ +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:\\*NON_LVALUE_EXPR \\\[len: 64\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: 128\\\]\\)" 1 "original" { target { c++ } } } } */ +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc cache \\(readonly:x\\\[0\\\] \\\[len: 128\\\]\\);$" 1 "original" } } */ + + + diff --git a/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90 b/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90 new file mode 100644 index 00000000000..069fec0a0d5 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90 @@ -0,0 +1,28 @@ +! { dg-additional-options "-fdump-tree-original" } + +subroutine foo (a, n) + integer :: n, a(:) + integer :: i, b(n) + !$acc parallel copyin(readonly: a(:), b(:n)) + do i = 1,32 + !$acc cache (readonly: a(:), b(:n)) + enddo + !$acc end parallel +end subroutine foo + +program main + integer :: i, n = 32, a(32) + integer :: b(32) + !$acc parallel copyin(readonly: a(:32), b(:n)) + do i = 1,32 + !$acc cache (readonly: a(:), b(:n)) + enddo + !$acc end parallel +end program main + +! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\) .+ map\\(readonly,to:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:a\\\[\\(\\(integer\\(kind=8\\)\\) parm.*data - \\(integer\\(kind=8\\)\\) &a\\) / 4\\\] \\\[len: .+\\\]\\) .+ map\\(readonly,to:b\\\[\\(\\(integer\\(kind=8\\)\\) parm.*data - \\(integer\\(kind=8\\)\\) &b\\) / 4\\\] \\\[len: .+\\\]\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma acc cache \\(readonly:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\) \\(readonly:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\);" 2 "original" } } + + + diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc index a743e3cdfd8..6a9812c2253 100644 --- a/gcc/tree-pretty-print.cc +++ b/gcc/tree-pretty-print.cc @@ -905,6 +905,8 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) case OMP_CLAUSE_MAP: pp_string (pp, "map("); + if (OMP_CLAUSE_MAP_READONLY (clause)) + pp_string (pp, "readonly,"); switch (OMP_CLAUSE_MAP_KIND (clause)) { case GOMP_MAP_ALLOC: @@ -1075,6 +1077,8 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) case OMP_CLAUSE__CACHE_: pp_string (pp, "("); + if (OMP_CLAUSE__CACHE__READONLY (clause)) + pp_string (pp, "readonly:"); dump_generic_node (pp, OMP_CLAUSE_DECL (clause), spc, flags, false); goto print_clause_size; diff --git a/gcc/tree.h b/gcc/tree.h index 3eebf5709b7..a79260e48eb 100644 --- a/gcc/tree.h +++ b/gcc/tree.h @@ -1813,6 +1813,14 @@ class auto_suppress_location_wrappers #define OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE(NODE) \ (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->base.addressable_flag) +/* Nonzero if OpenACC 'readonly' modifier set, used for 'copyin'. */ +#define OMP_CLAUSE_MAP_READONLY(NODE) \ + TREE_READONLY (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)) + +/* Same as above, for use in OpenACC cache directives. */ +#define OMP_CLAUSE__CACHE__READONLY(NODE) \ + TREE_READONLY (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE__CACHE_)) + /* True on an OMP_CLAUSE_USE_DEVICE_PTR with an OpenACC 'if_present' clause. */ #define OMP_CLAUSE_USE_DEVICE_PTR_IF_PRESENT(NODE) \