From patchwork Mon Aug 7 13:58:27 2023 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Chung-Lin Tang X-Patchwork-Id: 1817991 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=wD9G/8Be; 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 4RKHx51PCJz1yVt for ; Mon, 7 Aug 2023 23:59:04 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 179E638582A1 for ; Mon, 7 Aug 2023 13:59:02 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 179E638582A1 DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gcc.gnu.org; s=default; t=1691416742; bh=B5S9YkwaztrvC+NPB1/+x4ZaFlsyqAGO/68YX8g9ggQ=; h=Date:Subject:To:Cc:References:In-Reply-To:List-Id: List-Unsubscribe:List-Archive:List-Post:List-Help:List-Subscribe: From:Reply-To:From; b=wD9G/8BeOLRS+ZlTgQZcFOKkVol5g03Zh1S3i6G5J8P5/TIpZN8inDqWux4mrf9Mh xCu56tseLRwEkqHswyBrnecYcvyXmkFRrRRqK+ksoSlq0ba9+FKwpmWlxdPeDFeIJG z86LyW/QkQulaV3/vagvlG9MvLvJNvZrsJFeZzIE= X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from APC01-SG2-obe.outbound.protection.outlook.com (mail-sgaapc01on2050.outbound.protection.outlook.com [40.107.215.50]) by sourceware.org (Postfix) with ESMTPS id 4277C3858D28 for ; Mon, 7 Aug 2023 13:58:38 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 4277C3858D28 ARC-Seal: i=1; a=rsa-sha256; s=arcselector9901; d=microsoft.com; cv=none; b=Q2NUqH4aB2zZT3SgpTxrs6z89uo/bgSFZnLQuH0wE773of44x8io83zNZzb2xh529AMQpRTJeERLNYTqXbvxuiEa2pV6xJ5qlu0uDpTSdEN2Vgyl//JxkJVSq9tyACpAYZjXug9gN64B5GXq+mlr5CroZ4KNn6WKLWYZTMs3ZZ1UIS9If8ggy5DNSZwBcZv6FtjISIX2e3t1TYoRUsRDEaxocEUdga73EzhO0VRPmcg6NWCGcvAsxCfwnxoCOv/4kbuwQhWIzYuzbWj5caR6/2dFx1e8uS//QEX+wlswYrzs6MBDziXsGO4lnumKnFsMpA99mzKAY/pUUxSGJarMWA== 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=B5S9YkwaztrvC+NPB1/+x4ZaFlsyqAGO/68YX8g9ggQ=; b=cKd3Pnr1/00c4Xag99EAbDuvt6hRNmjncFDz8Kw4kodL+ncydEBnqY2Fhn0IeDqvh2UQH0GE00AA619qYLwgQ9c5rgBAu5DJirGHIibNND7j+viFudeOIldK5E1VbE2PqdvNgQanlaSRS8o3CHfDVdkWa9vYmaeQvjkLdK/sE5kfKmiXl7MSJ17sHlEEOffoZ5GkyzrK8uIdg5JG+1uZl+h2pgb1qlYCWlTspafhGhfWbsFwkEEPXnNEvjlOck8dwKDwEmWJYk2NJcUbJdLszinb60YFeWSyiNUJaO4E8Di/DXnFLCsinXsTz0AVU8g5qLzx5ouGFyg+LyiBeoUuLw== 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 KL1PR0601MB4049.apcprd06.prod.outlook.com (2603:1096:820:22::13) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_256_GCM_SHA384) id 15.20.6652.26; Mon, 7 Aug 2023 13:58:32 +0000 Received: from SG2PR06MB5430.apcprd06.prod.outlook.com ([fe80::e720:fef7:1199:a273]) by SG2PR06MB5430.apcprd06.prod.outlook.com ([fe80::e720:fef7:1199:a273%7]) with mapi id 15.20.6652.022; Mon, 7 Aug 2023 13:58:32 +0000 Message-ID: Date: Mon, 7 Aug 2023 21:58:27 +0800 User-Agent: Mozilla/5.0 (Macintosh; Intel Mac OS X 10.15; rv:102.0) Gecko/20100101 Thunderbird/102.13.0 Subject: [PATCH, OpenACC 2.7, v2] readonly modifier support in front-ends To: Tobias Burnus , Thomas Schwinge , Chung-Lin Tang Cc: gcc-patches@gcc.gnu.org, Catherine Moore References: <87lefaaesb.fsf@euler.schwinge.homeip.net> Content-Language: en-US In-Reply-To: X-ClientProxiedBy: TYCP301CA0028.JPNP301.PROD.OUTLOOK.COM (2603:1096:400:381::11) To SG2PR06MB5430.apcprd06.prod.outlook.com (2603:1096:4:1ba::14) MIME-Version: 1.0 X-MS-PublicTrafficType: Email X-MS-TrafficTypeDiagnostic: SG2PR06MB5430:EE_|KL1PR0601MB4049:EE_ X-MS-Office365-Filtering-Correlation-Id: 778f2ac1-650c-454b-6eda-08db974e6267 X-MS-Exchange-SenderADCheck: 1 X-MS-Exchange-AntiSpam-Relay: 0 X-Microsoft-Antispam: BCL:0; X-Microsoft-Antispam-Message-Info: KDYiBb2Rx3+fwJxuS2l5gf2u5HUKm4ib0cGR9Kqw82HbLdhK5UNVbrupTgQtD5k2p3+3ZjPyMNgEO30AVyKf6p/vNvys5ZmD+Umm/V51f+U8PuVJyGwN6+g0IVPEXcvmzntimVZzyWDgX2RzKrMzQLm/53JqvoFNN0Iu5pPXVTXjAIB1gjQXOSTX/Fc4t3mXDhse1pOaw1rHlWuy5hErSwd8Rfoe12FWC4sLDkalcMObfCujBxtXge+LAuCqEt5pcF0lJeQ+OnyXYhLuM8YmyP9Qumzd3gYSb1Ioe4Qt1vsoufjUIqmn6YhzkM8YT/CsuqP6EHEx3YQS0MRYZYlXNzj6UbdFTuurIC2vbFdqESB+geJLvlkOdSDq2a7PnLSVrBa0TqYPJpaFGrxhYq8zviGDbL54to1+SFXB4gGEzdaRLOG3ItfzZg3xIyGnOT8iYz2yedsjf9Sx89OHo7p4yfgeyFm3fjje9l9Z3hdMoSgrqpAdenHwq4V8o9scyW7KHXbsmCrkDXGvWwPBnu4sLeMN9EFLREMGvHxiSKnQ8+fhpJUe/gUfqXOyjBdnCD5V93f34+I0ARJcha+vVibj6me/bXe97uBmfYerJrPK6/RbsSz4Uiu939L+XQJI+RSGwcNZtSuo5DAlm3pJ4PQuKw== 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)(376002)(366004)(346002)(396003)(136003)(39860400002)(451199021)(186006)(1800799003)(83380400001)(84970400001)(31686004)(2616005)(110136005)(4326008)(38100700002)(2906002)(5660300002)(316002)(235185007)(8676002)(8936002)(66946007)(66556008)(66476007)(6512007)(86362001)(478600001)(31696002)(82960400001)(6666004)(33964004)(6486002)(41300700001)(36756003)(6506007)(26005)(53546011)(45980500001)(43740500002); DIR:OUT; SFP:1101; X-MS-Exchange-AntiSpam-MessageData-ChunkCount: 1 X-MS-Exchange-AntiSpam-MessageData-0: =?utf-8?q?rMaHkrQVMie5DH2QCC4XNa6x6jVt?= =?utf-8?q?tDDuyIFBFEfdyLERTLaQ1G+XzaPjwWr2522fiMQI5Rfrs2ms9F18d5X+npgs2oZaH?= =?utf-8?q?WF8hD4mI9SWbzZaA8h+UC/dPBgOby2arz4aMxxQQSXcLoDVvckkZHpIcRE72KAU5A?= =?utf-8?q?lhZWMo7IyPegsD6pkS9kAF7hLaS5Qipur5eo13xHg1hyObVxuQ4pH8wRvCBM9wgRw?= =?utf-8?q?4DmJOrD3ka0JmdJS1FxQM6zkc28z54WSfO/7YHTTEDvWrHd+NzANH4F8XAo66x/u3?= =?utf-8?q?N3H+GSoLQUi3JDO/Mv29ubSXDsv0YCE93GjmY+8+SP0o56QBtiwFyDcj4u+nus+LX?= =?utf-8?q?wqizAVA7SkIZTCKcm8/6S2lm30jd8+nIwLa2AKscjrf8khjTOE54+a5DnrH3+RXJX?= =?utf-8?q?8rAVFQ41PRs29zxCG5QpJiAOva4QMS+nIfoSIGlYVOLUJHQq2yRauBhCe0f+T4+IN?= =?utf-8?q?fm6K8fX4O3E+zTVov5QJku8oL1jGOrPIeM0xvGdtM61d/IaKfskLQciqpYnqRE4Gz?= =?utf-8?q?l3KsYyaX9Pj1BCPmGvGZWIpxr1OgK07zEjNIBniHepKDwBIGkmDXIPwWnQ2RFqC4G?= =?utf-8?q?xWZ5fMTa8cosg16+Rq3dJeG1GRiBOq5SvUo0lM0RgMscBY1YVLOSaRsDeTY87HXnK?= =?utf-8?q?1SYqN4rttXRHeA40cWve1xsgL2005TBJhWbVOxR/0BYZQBoU28QCcvmcOiZdmDQiJ?= =?utf-8?q?cPbxjc8RVf+guATUbSCH0TaB7LAzB1p3M/hO/3jI2G9SRTITMtdFeiP6sMaicMLYJ?= =?utf-8?q?RDiOwOXUZ8yIm6rK/cVZ+/Spdzr5r+/47Xmlg3m5X8XPGc21sNj0nyDeCNniGjFvj?= =?utf-8?q?oLOQ2tTyxnMVgNwo9sE9uiQrYAIGF3GA0K7VB6rrYQJlO7x1b7LXj/EnVm23BweFO?= =?utf-8?q?eS+MXmlx4NbYbZ02o8Y3fiIGeZ3eayIdrhgF6Sa0xE4gef19HI/ubDxKuvgfOaOJj?= =?utf-8?q?Mfb3XyY5jo+U8IFX6gonzMOWUDLo1iBg9EzevBIneMbxFhxozRUFqg5HDk1ST9lis?= =?utf-8?q?GXICf4DJpC7t2aW+uGrCSmiVfY6oPUZfHpEwd1twWganwz2kkr3hel0nlyy+V0SxB?= =?utf-8?q?qVjok7x6elqmubaHZI/5geQAg9yIsa4D7LyvMZJv6lcqvD8JV5YnRKAJ0vTtvz980?= =?utf-8?q?AcJKkyExFtHrKnqTNp39SwjrBDDTWExqBr3ehyQYGdsc9p2Eu1LVjL32BTih4lFLK?= =?utf-8?q?xb3Vr4hEbGVvaKamCTsI617JnRg4mQf30wF13eCV+MKugrhY1+CXpa4A6iv2kJhEX?= =?utf-8?q?Y2cN2KqplNP6slL4Fkc8XeTsYbVtYsJVseee3Bc236Jw5XJR2nlVigbu1oP2p9tvE?= =?utf-8?q?Y4K8tNojBKr+VlFgOGAv2JhAyLfyYepIv656yTEFNoMVeReCrGpgW8emI8czKZdIv?= =?utf-8?q?b4cVnpBXd/eUBSBH+QutODdgrB/IS6RMP3u45HDoY9/AIxSKBJIyjp98lBc7eDOpd?= =?utf-8?q?Tarht2lLrKtllkmjfh0xBPek3oZzmLAr30+c8ua0auDIxmcKsfeua9sO+A9lrvbhC?= =?utf-8?q?TTg15DP0NwRAmdMVkqa5x5AgvpW2TLomQQ=3D=3D?= X-OriginatorOrg: siemens.com X-MS-Exchange-CrossTenant-Network-Message-Id: 778f2ac1-650c-454b-6eda-08db974e6267 X-MS-Exchange-CrossTenant-AuthSource: SG2PR06MB5430.apcprd06.prod.outlook.com X-MS-Exchange-CrossTenant-AuthAs: Internal X-MS-Exchange-CrossTenant-OriginalArrivalTime: 07 Aug 2023 13:58:32.0142 (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: bTsno343FzOcWU5uvB2SG4yiaDEq9lyc6F2fGw6LxIqAhB3tqrXJR7sEnVHM0ywOMrt7lzPL3q26jOJXSmLvEpzn5dcpQGCi4mw/faLfw+U= X-MS-Exchange-Transport-CrossTenantHeadersStamped: KL1PR0601MB4049 X-Spam-Status: No, score=-10.1 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 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: Chung-Lin Tang Errors-To: gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org Sender: "Gcc-patches" Hi Thomas, Tobias, here's the updated v2 of the readonly modifier front-end patch. On 2023/7/20 11:08 PM, Tobias Burnus wrote: >>> +++ 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) >>> ... >> Instead of doing this in 'c_parser_omp_var_list_parens', I think it's >> clearer to have this special 'readonly :' parsing logic in the two places >> where it's used. > I concur. The same issue also occurred for OpenMP's > c_parser_omp_clause_to, and c_parser_omp_clause_from and the 'present' > modifier. For it, I created a combined function but the main reason for > that is that OpenMP also permits more modifiers (like 'iterators'), > which would cause more duplication of code ('iterator' is not yet > supported). > > For something as simple to parse as this modifier, I would just do it at > the two places – as Thomas suggested. Okay, I've changed the C/C++ parser parts to have the parsing logic directly added. >>> +++ 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 >>> { >> [...] Thus, the above looks good to me. > I concur but I wonder whether it would be cleaner to name the struct; > this makes it also more obvious what belongs together in the union. > > Namely, naming the struct 'map' and then changing the 45 users from > 'u.map_op' to 'u.map.op' and the new 'u.readonly' to 'u.map.readonly'. – > this seems to be cleaner. I've adjusted 'u.map' to be a named struct now, and updated the references. >> + if (gfc_match ("readonly :") == MATCH_YES) >> I note this one does not have a space after ':' in 'gfc_match', but the >> one above in 'gfc_match_omp_clauses' does. I don't know off-hand if that >> makes a difference in parsing -- probably not, as all of >> 'gcc/fortran/openmp.cc' generally doesn't seem to be very consistent >> about these two variants? > It *does* make a difference. And for obvious reasons. You don't want to permit: > > !$acc kernels asnyccopy(a) > > but require at least one space (or comma) between "async" and "copy".. > (In fixed form Fortran, it would be fine - as would be "!$acc k e nelsasy nc co p y(a)".) > > A " " matches zero or more whitespaces, but with gfc_match_space you can find out > whether there was whitespace or not. Okay, made sure both are 'gfc_match ("readonly : ")'. Thanks for catching that, didn't realize that space was significant. >>> +++ 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_)) >> I'm not sure if these special accessor functions are actually useful, or >> we should just directly use 'TREE_READONLY' instead? We're only using >> them in contexts where it's clear that the 'OMP_CLAUSE_SUBCODE_CHECK' is >> satisfied, for example. > I find directly using TREE_READONLY confusing. FWIW, I've changed to use TREE_NOTHROW instead, if it can give a better sense of safety :P I think there's a misunderstanding here anyways: we are not relying on a DECL marked TREE_READONLY here. We merely need the OMP_CLAUSE_MAP to be marked as OMP_CLAUSE_MAP_READONLY == 1. The other points-to patch then (also in front-ends) take the OMP_CLAUSE_MAP_READONLY to mark the clauses of "base-pointers of array-sections" as OMP_CLAUSE_MAP_POINTS_TO_READONLY, and later this gradually gets relayed to alias oracle routines in tree-ssa-alias.cc Re-tested this v2 patch on powerpc64le-linux/nvptx. Okay for trunk? Thanks, Chung-Lin 2023-08-07 Chung-Lin Tang gcc/c/ChangeLog: * c-parser.cc (c_parser_oacc_data_clause): Add parsing support for 'readonly' modifier, set OMP_CLAUSE_MAP_READONLY if readonly modifier found, update comments. (c_parser_oacc_cache): Add parsing support for 'readonly' modifier, set OMP_CLAUSE__CACHE__READONLY if readonly modifier found, update comments. gcc/cp/ChangeLog: * parser.cc (cp_parser_oacc_data_clause): Add parsing support for 'readonly' modifier, set OMP_CLAUSE_MAP_READONLY if readonly modifier found, update comments. (cp_parser_oacc_cache): Add parsing support for 'readonly' modifier, set OMP_CLAUSE__CACHE__READONLY if readonly modifier found, update comments. gcc/fortran/ChangeLog: * dump-parse-tree.cc (show_omp_namelist): Print "readonly," for OMP_LIST_MAP and OMP_LIST_CACHE if n->u.map.readonly is set. Adjust 'n->u.map_op' to 'n->u.map.op'. * gfortran.h (typedef struct gfc_omp_namelist): Adjust map_op as 'ENUM_BITFIELD (gfc_omp_map_op) op:8', add 'bool readonly' field, change to named struct field 'map'. * openmp.cc (gfc_match_omp_map_clause): Add 'bool readonly = false' parameter, set n->u.map.readonly field. Adjust 'n->u.map_op' to 'n->u.map.op'. (gfc_match_omp_clause_reduction): Adjust 'n->u.map_op' to 'n->u.map.op'. (gfc_match_omp_clauses): Add readonly modifier parsing for OpenACC copyin clause, adjust call to gfc_match_omp_map_clause. Adjust 'n->u.map_op' to 'n->u.map.op'. (gfc_match_oacc_declare): Adjust 'n->u.map_op' to 'n->u.map.op'. (gfc_match_oacc_cache): Add readonly modifier parsing for OpenACC cache directive. (resolve_omp_clauses): Adjust 'n->u.map_op' to 'n->u.map.op'. * trans-decl.cc (add_clause): Adjust 'n->u.map_op' to 'n->u.map.op'. (finish_oacc_declare): Likewise. * trans-openmp.cc (gfc_trans_omp_clauses): Set OMP_CLAUSE_MAP_READONLY, OMP_CLAUSE__CACHE__READONLY to 1 when readonly is set. Adjust 'n->u.map_op' to 'n->u.map.op'. (gfc_add_clause_implicitly): Adjust 'n->u.map_op' to 'n->u.map.op'. 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 24a6eb6e459..5779f499ae1 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -14084,7 +14084,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 +14139,36 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind, default: gcc_unreachable (); } - tree nl, c; - nl = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_MAP, list, true); - for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c)) - OMP_CLAUSE_SET_MAP_KIND (c, kind); + tree nl = list; + bool readonly = false; + matching_parens parens; + if (parens.require_open (parser)) + { + /* Turn on readonly modifier parsing for copyin clause. */ + if (c_kind == PRAGMA_OACC_CLAUSE_COPYIN) + { + 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; + } + } + location_t loc = c_parser_peek_token (parser)->location; + nl = c_parser_omp_variable_list (parser, loc, OMP_CLAUSE_MAP, list, true); + parens.skip_until_found_close (parser); + } + + for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c)) + { + OMP_CLAUSE_SET_MAP_KIND (c, kind); + if (readonly) + OMP_CLAUSE_MAP_READONLY (c) = 1; + } return nl; } @@ -18161,15 +18190,40 @@ 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. */ static tree c_parser_oacc_cache (location_t loc, c_parser *parser) { - tree stmt, clauses; + tree stmt, clauses = NULL_TREE; + bool readonly = false; + matching_parens parens; + + if (parens.require_open (parser)) + { + 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; + } + location_t loc = c_parser_peek_token (parser)->location; + clauses = c_parser_omp_variable_list (parser, loc, OMP_CLAUSE__CACHE_, + NULL_TREE); + parens.skip_until_found_close (parser); + } + + 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 d7ef5b34d42..ac8a656874a 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -37750,7 +37750,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, @@ -37801,11 +37805,33 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind, default: gcc_unreachable (); } - tree nl, c; - nl = cp_parser_omp_var_list (parser, OMP_CLAUSE_MAP, list, true); + tree nl = list; + bool readonly = false; + if (cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN)) + { + /* Turn on readonly modifier parsing for copyin clause. */ + if (c_kind == PRAGMA_OACC_CLAUSE_COPYIN) + { + cp_token *token = cp_lexer_peek_token (parser->lexer); + if (token->type == CPP_NAME + && !strcmp (IDENTIFIER_POINTER (token->u.value), "readonly") + && cp_lexer_peek_nth_token (parser->lexer, 2)->type == CPP_COLON) + { + cp_lexer_consume_token (parser->lexer); + cp_lexer_consume_token (parser->lexer); + readonly = true; + } + } + nl = cp_parser_omp_var_list_no_open (parser, OMP_CLAUSE_MAP, list, NULL, + true); + } - for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c)) - OMP_CLAUSE_SET_MAP_KIND (c, kind); + for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c)) + { + OMP_CLAUSE_SET_MAP_KIND (c, kind); + if (readonly) + OMP_CLAUSE_MAP_READONLY (c) = 1; + } return nl; } @@ -45825,6 +45851,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 @@ -45834,9 +45863,28 @@ cp_parser_oacc_cache (cp_parser *parser, cp_token *pragma_tok) clauses. */ auto_suppress_location_wrappers sentinel; - tree stmt, clauses; + tree stmt, clauses = NULL_TREE; + bool readonly = false; + + if (cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN)) + { + cp_token *token = cp_lexer_peek_token (parser->lexer); + if (token->type == CPP_NAME + && !strcmp (IDENTIFIER_POINTER (token->u.value), "readonly") + && cp_lexer_peek_nth_token (parser->lexer, 2)->type == CPP_COLON) + { + cp_lexer_consume_token (parser->lexer); + cp_lexer_consume_token (parser->lexer); + readonly = true; + } + clauses = cp_parser_omp_var_list_no_open (parser, OMP_CLAUSE__CACHE_, + NULL, NULL); + } + + 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/dump-parse-tree.cc b/gcc/fortran/dump-parse-tree.cc index 68122e3e6fd..0e888fafe7b 100644 --- a/gcc/fortran/dump-parse-tree.cc +++ b/gcc/fortran/dump-parse-tree.cc @@ -1398,6 +1398,9 @@ show_omp_namelist (int list_type, gfc_omp_namelist *n) fputs (") ALLOCATE(", dumpfile); continue; } + if ((list_type == OMP_LIST_MAP || list_type == OMP_LIST_CACHE) + && n->u.map.readonly) + fputs ("readonly,", dumpfile); if (list_type == OMP_LIST_REDUCTION) switch (n->u.reduction_op) { @@ -1465,7 +1468,7 @@ show_omp_namelist (int list_type, gfc_omp_namelist *n) default: break; } else if (list_type == OMP_LIST_MAP) - switch (n->u.map_op) + switch (n->u.map.op) { case OMP_MAP_ALLOC: fputs ("alloc:", dumpfile); break; case OMP_MAP_TO: fputs ("to:", dumpfile); break; diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h index 9a00e6dea6f..a8667a6e6d3 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) op:8; + bool readonly; + } map; gfc_expr *align; struct { diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc index 2952cd300ac..af769d9efbd 100644 --- a/gcc/fortran/openmp.cc +++ b/gcc/fortran/openmp.cc @@ -1197,7 +1197,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, @@ -1206,7 +1206,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.map.readonly = readonly; + } return true; } @@ -1520,7 +1523,7 @@ gfc_match_omp_clause_reduction (char pc, gfc_omp_clauses *c, bool openacc, gfc_omp_namelist *p = gfc_get_omp_namelist (), **tl; p->sym = n->sym; p->where = p->where; - p->u.map_op = OMP_MAP_ALWAYS_TOFROM; + p->u.map.op = OMP_MAP_ALWAYS_TOFROM; tl = &c->lists[OMP_LIST_MAP]; while (*tl) @@ -2180,11 +2183,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], @@ -3101,7 +3109,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, { gfc_omp_namelist *n; for (n = *head; n; n = n->next) - n->u.map_op = map_op; + n->u.map.op = map_op; continue; } gfc_current_locus = old_loc; @@ -3942,7 +3950,7 @@ gfc_match_oacc_declare (void) if (gfc_current_ns->proc_name && gfc_current_ns->proc_name->attr.flavor == FL_MODULE) { - if (n->u.map_op != OMP_MAP_ALLOC && n->u.map_op != OMP_MAP_TO) + if (n->u.map.op != OMP_MAP_ALLOC && n->u.map.op != OMP_MAP_TO) { gfc_error ("Invalid clause in module with !$ACC DECLARE at %L", &where); @@ -3976,7 +3984,7 @@ gfc_match_oacc_declare (void) return MATCH_ERROR; } - switch (n->u.map_op) + switch (n->u.map.op) { case OMP_MAP_FORCE_ALLOC: case OMP_MAP_ALLOC: @@ -4091,20 +4099,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.map.readonly = true; + if (gfc_current_state() != COMP_DO && gfc_current_state() != COMP_DO_CONCURRENT) { @@ -8142,8 +8165,8 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, } if (openacc && list == OMP_LIST_MAP - && (n->u.map_op == OMP_MAP_ATTACH - || n->u.map_op == OMP_MAP_DETACH)) + && (n->u.map.op == OMP_MAP_ATTACH + || n->u.map.op == OMP_MAP_DETACH)) { symbol_attribute attr; if (n->expr) @@ -8153,7 +8176,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, if (!attr.pointer && !attr.allocatable) gfc_error ("%qs clause argument must be ALLOCATABLE or " "a POINTER at %L", - (n->u.map_op == OMP_MAP_ATTACH) ? "attach" + (n->u.map.op == OMP_MAP_ATTACH) ? "attach" : "detach", &n->where); } if (lastref @@ -8224,7 +8247,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, else if (openacc) { if (list == OMP_LIST_MAP - && n->u.map_op == OMP_MAP_FORCE_DEVICEPTR) + && n->u.map.op == OMP_MAP_FORCE_DEVICEPTR) resolve_oacc_deviceptr_clause (n->sym, n->where, name); else resolve_oacc_data_clauses (n->sym, n->where, name); @@ -8246,7 +8269,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, { case EXEC_OMP_TARGET: case EXEC_OMP_TARGET_DATA: - switch (n->u.map_op) + switch (n->u.map.op) { case OMP_MAP_TO: case OMP_MAP_ALWAYS_TO: @@ -8273,7 +8296,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, } break; case EXEC_OMP_TARGET_ENTER_DATA: - switch (n->u.map_op) + switch (n->u.map.op) { case OMP_MAP_TO: case OMP_MAP_ALWAYS_TO: @@ -8283,16 +8306,16 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, case OMP_MAP_PRESENT_ALLOC: break; case OMP_MAP_TOFROM: - n->u.map_op = OMP_MAP_TO; + n->u.map.op = OMP_MAP_TO; break; case OMP_MAP_ALWAYS_TOFROM: - n->u.map_op = OMP_MAP_ALWAYS_TO; + n->u.map.op = OMP_MAP_ALWAYS_TO; break; case OMP_MAP_PRESENT_TOFROM: - n->u.map_op = OMP_MAP_PRESENT_TO; + n->u.map.op = OMP_MAP_PRESENT_TO; break; case OMP_MAP_ALWAYS_PRESENT_TOFROM: - n->u.map_op = OMP_MAP_ALWAYS_PRESENT_TO; + n->u.map.op = OMP_MAP_ALWAYS_PRESENT_TO; break; default: gfc_error ("TARGET ENTER DATA with map-type other " @@ -8302,7 +8325,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, } break; case EXEC_OMP_TARGET_EXIT_DATA: - switch (n->u.map_op) + switch (n->u.map.op) { case OMP_MAP_FROM: case OMP_MAP_ALWAYS_FROM: @@ -8312,16 +8335,16 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, case OMP_MAP_DELETE: break; case OMP_MAP_TOFROM: - n->u.map_op = OMP_MAP_FROM; + n->u.map.op = OMP_MAP_FROM; break; case OMP_MAP_ALWAYS_TOFROM: - n->u.map_op = OMP_MAP_ALWAYS_FROM; + n->u.map.op = OMP_MAP_ALWAYS_FROM; break; case OMP_MAP_PRESENT_TOFROM: - n->u.map_op = OMP_MAP_PRESENT_FROM; + n->u.map.op = OMP_MAP_PRESENT_FROM; break; case OMP_MAP_ALWAYS_PRESENT_TOFROM: - n->u.map_op = OMP_MAP_ALWAYS_PRESENT_FROM; + n->u.map.op = OMP_MAP_ALWAYS_PRESENT_FROM; break; default: gfc_error ("TARGET EXIT DATA with map-type other " diff --git a/gcc/fortran/trans-decl.cc b/gcc/fortran/trans-decl.cc index b0fd25e92a3..1ff1dda026a 100644 --- a/gcc/fortran/trans-decl.cc +++ b/gcc/fortran/trans-decl.cc @@ -6614,7 +6614,7 @@ add_clause (gfc_symbol *sym, gfc_omp_map_op map_op) n = gfc_get_omp_namelist (); n->sym = sym; - n->u.map_op = map_op; + n->u.map.op = map_op; if (!module_oacc_clauses) module_oacc_clauses = gfc_get_omp_clauses (); @@ -6716,10 +6716,10 @@ finish_oacc_declare (gfc_namespace *ns, gfc_symbol *sym, bool block) for (n = omp_clauses->lists[OMP_LIST_MAP]; n; n = n->next) { - switch (n->u.map_op) + switch (n->u.map.op) { case OMP_MAP_DEVICE_RESIDENT: - n->u.map_op = OMP_MAP_FORCE_ALLOC; + n->u.map.op = OMP_MAP_FORCE_ALLOC; break; default: diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc index cf741cebf91..a4628e460bd 100644 --- a/gcc/fortran/trans-openmp.cc +++ b/gcc/fortran/trans-openmp.cc @@ -3067,7 +3067,10 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, || (n->expr && gfc_expr_attr (n->expr).pointer))) always_modifier = true; - switch (n->u.map_op) + if (n->u.map.readonly) + OMP_CLAUSE_MAP_READONLY (node) = 1; + + switch (n->u.map.op) { case OMP_MAP_ALLOC: OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALLOC); @@ -3194,8 +3197,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, && n->sym->attr.omp_declare_target && (always_modifier || n->sym->attr.pointer) && op != EXEC_OMP_TARGET_EXIT_DATA - && n->u.map_op != OMP_MAP_DELETE - && n->u.map_op != OMP_MAP_RELEASE) + && n->u.map.op != OMP_MAP_DELETE + && n->u.map.op != OMP_MAP_RELEASE) { gcc_assert (n->sym->ts.u.cl->backend_decl); node5 = build_omp_clause (input_location, OMP_CLAUSE_MAP); @@ -3261,7 +3264,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, { enum gomp_map_kind gmk = GOMP_MAP_POINTER; if (op == EXEC_OMP_TARGET_EXIT_DATA - && n->u.map_op == OMP_MAP_DELETE) + && n->u.map.op == OMP_MAP_DELETE) gmk = GOMP_MAP_DELETE; else if (op == EXEC_OMP_TARGET_EXIT_DATA) gmk = GOMP_MAP_RELEASE; @@ -3284,7 +3287,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, { enum gomp_map_kind gmk; if (op == EXEC_OMP_TARGET_EXIT_DATA - && n->u.map_op == OMP_MAP_DELETE) + && n->u.map.op == OMP_MAP_DELETE) gmk = GOMP_MAP_DELETE; else if (op == EXEC_OMP_TARGET_EXIT_DATA) gmk = GOMP_MAP_RELEASE; @@ -3316,18 +3319,18 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, node2 = build_omp_clause (input_location, OMP_CLAUSE_MAP); OMP_CLAUSE_DECL (node2) = decl; OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type); - if (n->u.map_op == OMP_MAP_DELETE) + if (n->u.map.op == OMP_MAP_DELETE) map_kind = GOMP_MAP_DELETE; else if (op == EXEC_OMP_TARGET_EXIT_DATA - || n->u.map_op == OMP_MAP_RELEASE) + || n->u.map.op == OMP_MAP_RELEASE) map_kind = GOMP_MAP_RELEASE; else map_kind = GOMP_MAP_TO_PSET; OMP_CLAUSE_SET_MAP_KIND (node2, map_kind); if (op != EXEC_OMP_TARGET_EXIT_DATA - && n->u.map_op != OMP_MAP_DELETE - && n->u.map_op != OMP_MAP_RELEASE) + && n->u.map.op != OMP_MAP_DELETE + && n->u.map.op != OMP_MAP_RELEASE) { node3 = build_omp_clause (input_location, OMP_CLAUSE_MAP); @@ -3345,7 +3348,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, = gfc_conv_descriptor_data_get (decl); OMP_CLAUSE_SIZE (node3) = size_int (0); - if (n->u.map_op == OMP_MAP_ATTACH) + if (n->u.map.op == OMP_MAP_ATTACH) { /* Standalone attach clauses used with arrays with descriptors must copy the descriptor to the @@ -3361,7 +3364,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, node3 = NULL; goto finalize_map_clause; } - else if (n->u.map_op == OMP_MAP_DETACH) + else if (n->u.map.op == OMP_MAP_DETACH) { OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_DETACH); /* Similarly to above, we don't want to unmap PTR @@ -3553,8 +3556,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, to perform a single attach/detach operation, of the pointer itself, not of the pointed-to object. */ if (openacc - && (n->u.map_op == OMP_MAP_ATTACH - || n->u.map_op == OMP_MAP_DETACH)) + && (n->u.map.op == OMP_MAP_ATTACH + || n->u.map.op == OMP_MAP_DETACH)) { OMP_CLAUSE_DECL (node) = build_fold_addr_expr (OMP_CLAUSE_DECL (node)); @@ -3585,7 +3588,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, fold_convert (size_type_node, se.string_length), TYPE_SIZE_UNIT (tmp)); - if (n->u.map_op == OMP_MAP_DELETE) + if (n->u.map.op == OMP_MAP_DELETE) kind = GOMP_MAP_DELETE; else if (op == EXEC_OMP_TARGET_EXIT_DATA) kind = GOMP_MAP_RELEASE; @@ -3642,8 +3645,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, to perform a single attach/detach operation, of the pointer itself, not of the pointed-to object. */ if (openacc - && (n->u.map_op == OMP_MAP_ATTACH - || n->u.map_op == OMP_MAP_DETACH)) + && (n->u.map.op == OMP_MAP_ATTACH + || n->u.map.op == OMP_MAP_DETACH)) { OMP_CLAUSE_DECL (node) = build_fold_addr_expr (inner); @@ -3689,8 +3692,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, { /* Bare attach and detach clauses don't want any additional nodes. */ - if ((n->u.map_op == OMP_MAP_ATTACH - || n->u.map_op == OMP_MAP_DETACH) + if ((n->u.map.op == OMP_MAP_ATTACH + || n->u.map.op == OMP_MAP_DETACH) && (POINTER_TYPE_P (TREE_TYPE (inner)) || GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (inner)))) { @@ -3724,8 +3727,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, map_kind = ((GOMP_MAP_ALWAYS_P (map_kind) || gfc_expr_attr (n->expr).pointer) ? GOMP_MAP_ALWAYS_TO : GOMP_MAP_TO); - else if (n->u.map_op == OMP_MAP_RELEASE - || n->u.map_op == OMP_MAP_DELETE) + else if (n->u.map.op == OMP_MAP_RELEASE + || n->u.map.op == OMP_MAP_DELETE) ; else if (op == EXEC_OMP_TARGET_EXIT_DATA) map_kind = GOMP_MAP_RELEASE; @@ -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.map.readonly) + OMP_CLAUSE__CACHE__READONLY (node) = 1; omp_clauses = gfc_trans_add_clause (node, omp_clauses); } break; @@ -6333,7 +6338,7 @@ gfc_add_clause_implicitly (gfc_omp_clauses *clauses_out, n2->where = n->where; n2->sym = n->sym; if (is_target) - n2->u.map_op = OMP_MAP_TOFROM; + n2->u.map.op = OMP_MAP_TOFROM; if (tail) { tail->next = n2; 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 25d191b10fd..9604c3eecc5 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 4c04245e2b1..1301491587f 100644 --- a/gcc/tree.h +++ b/gcc/tree.h @@ -1811,6 +1811,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_NOTHROW (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)) + +/* Same as above, for use in OpenACC cache directives. */ +#define OMP_CLAUSE__CACHE__READONLY(NODE) \ + TREE_NOTHROW (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) \