From patchwork Tue Jul 25 15:52:06 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: 1812723 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=jJE51Js7; 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 4R9M4H3Qf5z1ydq for ; Wed, 26 Jul 2023 01:52:45 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 5DEC43856DFE for ; Tue, 25 Jul 2023 15:52:43 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 5DEC43856DFE DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gcc.gnu.org; s=default; t=1690300363; bh=3ndhBhp7vaT6LRZMYD9av5KBhIqSF3N28tdToRZPAek=; h=Date:Subject:To:References:In-Reply-To:List-Id:List-Unsubscribe: List-Archive:List-Post:List-Help:List-Subscribe:From:Reply-To:Cc: From; b=jJE51Js7k3cb7goobxFQX41ly3tW0D2i9lrGjUTGMP3b7A0938qq7sWJTGwNeqnX3 HNFN22TPrwmEDT4MgQ2GJHnnknxJO1glO11mBB58IGxZObDOsFNQHYb/UvE1mKSJQ3 O5k9XNruXizICzQW/KZWZvM8tKRsyLQvUOJKYGFo= X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from APC01-TYZ-obe.outbound.protection.outlook.com (mail-tyzapc01on2081.outbound.protection.outlook.com [40.107.117.81]) by sourceware.org (Postfix) with ESMTPS id E3F703858280 for ; Tue, 25 Jul 2023 15:52:16 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org E3F703858280 ARC-Seal: i=1; a=rsa-sha256; s=arcselector9901; d=microsoft.com; cv=none; b=a8IJBodTN0egNJH1PViZ5p8i2vAC6lD9tlVfvMgJr5RG+s7xpI1ZX0QQA1xJcg2/trhNxuX96TV5WJK+XLGUWMZveobYxcO5dhEHx0xwF1YCs2GX8ATPqg/WzCFeljspmQaqYP8rvruz9cazNAL/i3tvmi6osPWTwdjAf+xqfF1EqZO2Ic1m/8dLASL7lhGh3mOkhK+eoslSgE+rsdf/sItpgo/0JEKytBifD4fN9phzSUwM1uFZRcW5ROFyGjEfChzTnza8jDCf1/a9cmIVcQeF0dh8P3k3owohCpCMhI1AESb/KtndkEXrktE4URpteWxgnwt7pquMhPLLpJBpGw== 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=3ndhBhp7vaT6LRZMYD9av5KBhIqSF3N28tdToRZPAek=; b=KZKVPNGql9dzEF0V7zx4pFtFASlLXthmByXtMFvJVfpDgoJEYUGqRvEnwp9J+BvOpzuif9Jn+wsZ3XrqNL+g4kg8G4qNFci2Ttgi7OnLSRySu0suF5gJn8YF/QcMWGW/o7eSyLS/Mzf2RiggkuHQor7LKGkciXBLdMcWYvmCVZCVL7OGljdMDBRyf77V0y9CVwWT0Ikus/SkYYFLZHplzZZvqLNc7ckNxO861S4LWdFIbc667ILZjaXhsWfIEM+gtuQUfYUcx5R064BtG5T12SO+FDK1gCu2kuGpp3rCjX7qL3Uk5GW5TwU43Tjxl5FI/8gjHqlMk1BjRjB9dUNxvw== 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 SEYPR06MB6538.apcprd06.prod.outlook.com (2603:1096:101:177::8) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_256_GCM_SHA384) id 15.20.6609.33; Tue, 25 Jul 2023 15:52:11 +0000 Received: from SG2PR06MB5430.apcprd06.prod.outlook.com ([fe80::f7f2:5dc1:2d75:a61a]) by SG2PR06MB5430.apcprd06.prod.outlook.com ([fe80::f7f2:5dc1:2d75:a61a%7]) with mapi id 15.20.6609.032; Tue, 25 Jul 2023 15:52:11 +0000 Message-ID: <5196826c-e81a-ab5c-63e9-bd8509232da0@siemens.com> Date: Tue, 25 Jul 2023 23:52:06 +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] Connect readonly modifier to points-to analysis Content-Language: en-US To: cltang@codesourcery.com, gcc-patches , Thomas Schwinge , Catherine Moore , Tobias Burnus References: In-Reply-To: X-ClientProxiedBy: TYCPR01CA0075.jpnprd01.prod.outlook.com (2603:1096:405:3::15) To SG2PR06MB5430.apcprd06.prod.outlook.com (2603:1096:4:1ba::14) MIME-Version: 1.0 X-MS-PublicTrafficType: Email X-MS-TrafficTypeDiagnostic: SG2PR06MB5430:EE_|SEYPR06MB6538:EE_ X-MS-Office365-Filtering-Correlation-Id: b1d0285e-d64f-4305-fc92-08db8d271b43 X-MS-Exchange-SenderADCheck: 1 X-MS-Exchange-AntiSpam-Relay: 0 X-Microsoft-Antispam: BCL:0; X-Microsoft-Antispam-Message-Info: HpM9MBQ7/S70BnkMB7dlFJ+zoDLy9LpRNESDvSBUmV4OodAsZMgWfRmO164sHf3MWVNSg0I9VdLu4b/g+bKirgVUGsbEzW21yS5pdfqa/cnXRPrBHDmVoAbNLWZmxwxuZf859ORZhJHi7qB4Q6h4BD+Rt0fmpXuHbdm2+XAfJSvK73DgCXB1/dP5GL8fseAG4go4TICDjCMBm+D2k0HBzH3InJ8Q0gbLmnLjDpQPZaJwN2z+T3Sx6KGRjDj/T1d7kL6LAuzJeLMrnATDHY5StOllOeMiFe+UYj9afztmS6OppHv3fOwDyR/VvucijOEsJWVY3OdY94WFN3vpyow/yiOWNigxfaIVhxu8Lt/oS2TulZhUb0lBEfRq2EkoabKNlRHjubZXPTw95f+x1q91PEndPtGLtZnZzewdLeGnYKBHhDf1CgH1bSiywxkGsj6shdMESYRfIN4qZChlaRZRDF5s21aXTJJVN/qTqyc2dgHZOFnRdU/d77JGalCn1udMIUIboO0+NkyOYMCEBfk2DUncgBun+QNUVqs/CAis57SW1nRFm9Y25+5Ek0UkWGe0FKkSBTRRhOXDaq92vSoUiHDsSKL8Xju7OXPfxkoUV0V7XJBPv/JBCLdWx+GQSoEI2kNupEbKrS78vVSWgDqrVQ== 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)(346002)(39860400002)(136003)(366004)(376002)(396003)(451199021)(84970400001)(31686004)(110136005)(478600001)(186003)(6666004)(33964004)(6486002)(235185007)(5660300002)(8936002)(6506007)(36756003)(6512007)(2906002)(26005)(83380400001)(82960400001)(66476007)(38100700002)(41300700001)(8676002)(53546011)(31696002)(316002)(86362001)(2616005)(66946007)(66556008)(43740500002)(45980500001); DIR:OUT; SFP:1101; X-MS-Exchange-AntiSpam-MessageData-ChunkCount: 1 X-MS-Exchange-AntiSpam-MessageData-0: =?utf-8?q?QK26MF85CJW/p6uKMMMuwa7LET8V?= =?utf-8?q?XBlJZiufEXk1TZTo+A/dGfAIKDyehF4arxTnrTgctGrb2dzz336czV1PsE93a4wpz?= =?utf-8?q?VF3fj1NBPsZtLaGrj83NlFfuloG86i0+rmNrrFTaacB+3LHOo1zmSq5SkhgqQHxg3?= =?utf-8?q?Zd8+z3rtFpzqVb9FGHFcf8E1n537OKMQ7Vjt/dUil8IVzfZHYB9uCGHhJ9FRJCkhJ?= =?utf-8?q?Is+1bQ/G1JTT36RVQy3vN9Ef5b+xnmvWREldVfvh1GOjChNNTofdhUo6VoaihH7OM?= =?utf-8?q?Ay2uFedmO3ZEUmTAhNX8W6oxxdjxO+xNmcK9dB654b1h5bNQ8H+dTBfScQvDdhM/U?= =?utf-8?q?Z1hYmA+zPBfITJBdL9DHCBOlalI+RgVHKpLTvUnhD5QjX+5SbfYZCdU1tMUhgfWI9?= =?utf-8?q?+GZ1Xwin1/pff6giMh5sATstxvZR9WhLwOAXfVeinOhzuVPjA4pq6j85dEQYyfCYb?= =?utf-8?q?HOEQ9bP5Xu9djnC4BGVKHfIaWzdJyTyquaczRJgvreqPWRtE7Mlu2G/TgmfVZPjNd?= =?utf-8?q?jSGNsOhsr15HA1wIaj0Y1GesGKq/X2XZ3QjQx4ZBOreHuZTmEnrV1MHvO2WfCFt25?= =?utf-8?q?1zmSkY+cAaWB9Mw0HdLTNIGJ5xvq24qk+c7IhTAIkQmJ7WC+/sILCjzhfXGtFLZKz?= =?utf-8?q?IAwg8tqXNzr31wFkwWJkR7o+WmdVuCVt/p8G2AwcEnLVFx5S4rtUOSicRa40bRHye?= =?utf-8?q?DpVL6DMuIoCoZl3Ekt3/kO61CnDYFYYIvfgfMLcqybcPHSiARFOwS6c5hQ7bq0JCR?= =?utf-8?q?aGDiStufvhwgtSogptxKxvCJQzNGM7NsVbQCEACBPlplLrJryF5+464+MYfP22eKS?= =?utf-8?q?aPvonrleilzycnnbrcB3AyLp4ryh6DGd6XwjhkA/4FN8UbQ4BlwDz5LWC5706mH1m?= =?utf-8?q?nYQ7WBT0sXa0V4fKiu8yWjpbkOkztlc+TPZouSo41EGd4ZuxMP2yD+PdKvJO+DgTi?= =?utf-8?q?7S3ajKvjsel3RvCY+2kWnwijhNhbW5UoQtQLhXojV5OrZnyNknkhy1o1wPcFbENfy?= =?utf-8?q?OuyTXBmgZXQJ/IiVv6z1Ld+4NO/qlKUDf8U7A3oRHCISt5wRjPEAUoexZ4x7zUx9S?= =?utf-8?q?cO+458KqJ4C1n4nooon2VZcHt1+Ku1TMR/cUfxmevC5ecJkhBy8fvj5KSPEv3AUAN?= =?utf-8?q?hYHfj/xjZXfm99oeMmrEgyi/l0HXlQ0OtqfwcHAgf4QQOYppwAfxQIEuNv8i4rlrh?= =?utf-8?q?ZbSwkIGxdcUktA50xu+qZO9P+J7Uj3GwddsxUeZEN/Hdflo6j/A9BO59cywdD6CkL?= =?utf-8?q?RWObwVgdg7kz+zFKvXG7S3FddOjSZwCASt2pwXpKxZvi6Y2TuL8EMVULpox0M3g4n?= =?utf-8?q?3E4sCeCN3D7FHrbGRIhrgAWYHdUMhCNhBQeUBHIRYMzhjrwJ7J9/B5hdSWBmLsE8l?= =?utf-8?q?W7GZr52jvrkiQ4QLbJK4qKLl6ge1VflBQUPOIAv8204tNwf32f0HnwaUgqUmZz6vy?= =?utf-8?q?PCNlMBdUGPOOtwNMfqGLfK+mYfZEP3VxAMGNi5zyj2+6TOy88izzE00dAGULonl91?= =?utf-8?q?GBXflgNNW3fmJDdLzrEtBAVpPrUciOVclg=3D=3D?= X-OriginatorOrg: siemens.com X-MS-Exchange-CrossTenant-Network-Message-Id: b1d0285e-d64f-4305-fc92-08db8d271b43 X-MS-Exchange-CrossTenant-AuthSource: SG2PR06MB5430.apcprd06.prod.outlook.com X-MS-Exchange-CrossTenant-AuthAs: Internal X-MS-Exchange-CrossTenant-OriginalArrivalTime: 25 Jul 2023 15:52:10.6402 (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: qjZkTDG+ALzaiixSrGPiPovs9isXzMCHNC+HWe9DWIxLmqNruO45/Xfr811E5Rd/1z/OwySOPbXW6ZgcLi+IKMnUFJuDE6vOuJmpm/H2Nrg= X-MS-Exchange-Transport-CrossTenantHeadersStamped: SEYPR06MB6538 X-Spam-Status: No, score=-10.0 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" On 2023/7/11 2:33 AM, Chung-Lin Tang via Gcc-patches wrote: > 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. This second patch tries to link the readonly modifier to points-to analysis. There already exists SSA_NAME_POINTS_TO_READONLY_MEMORY and it's support in the alias oracle routines in tree-ssa-alias.cc, so basically what this patch does is try to make the variables holding the array section base pointers to have this flag set. There is an another OMP_CLAUSE_MAP_POINTS_TO_READONLY set by front-ends on the associated pointer clauses if OMP_CLAUSE_MAP_READONLY is set. Also a DECL_POINTS_TO_READONLY flag is set for VAR_DECLs when creating the tmp vars carrying these receiver references on the offloaded side. These eventually get translated to SSA_NAME_POINTS_TO_READONLY_MEMORY. This still doesn't always work as expected in terms of optimization: struct pointer fields and Fortran arrays (kind of like C structs) which have several accesses to create the pointer access on the receive/offloaded side, and SRA appears to not work on these sequences, so gets in the way of much redundancy elimination. Currently have one testcase where we can demonstrate 'readonly' can avoid a clobber by function call. Tested on powerpc64le-linux/nvptx. Note this patch is create a-top of the front-end patch. (will respond to the other front-end patch comments later) Thanks, Chung-Lin 2023-07-25 Chung-Lin Tang gcc/c/ChangeLog: * c-typeck.cc (handle_omp_array_sections): Set OMP_CLAUSE_MAP_POINTS_TO_READONLY on pointer clause. gcc/cp/ChangeLog: * semantics.cc (handle_omp_array_sections): Set OMP_CLAUSE_MAP_POINTS_TO_READONLY on pointer clause. gcc/fortran/ChangeLog: * trans-openmp.cc (gfc_trans_omp_array_section): Set OMP_CLAUSE_MAP_POINTS_TO_READONLY on pointer clause. gcc/ChangeLog: * gimple-expr.cc (copy_var_decl): Copy DECL_POINTS_TO_READONLY for VAR_DECLs. * gimplify.cc (struct gimplify_omp_ctx): Add 'hash_set *pt_readonly_ptrs' field. (internal_get_tmp_var): Set DECL_POINTS_TO_READONLY/SSA_NAME_POINTS_TO_READONLY_MEMORY for new temp vars. (build_omp_struct_comp_nodes): Set OMP_CLAUSE_MAP_POINTS_TO_READONLY on pointer clause. (gimplify_scan_omp_clauses): Collect OMP_CLAUSE_MAP_POINTS_TO_READONLY to ctx->pt_readonly_ptrs. * omp-low.cc (lower_omp_target): Set DECL_POINTS_TO_READONLY for variables of receiver refs. * tree-pretty-print.cc (dump_omp_clause): Print OMP_CLAUSE_MAP_POINTS_TO_READONLY. (dump_generic_node): Print SSA_NAME_POINTS_TO_READONLY_MEMORY. * tree.h (DECL_POINTS_TO_READONLY): New macro. (OMP_CLAUSE_MAP_POINTS_TO_READONLY): New macro. gcc/testsuite/ChangeLog: * c-c++-common/goacc/readonly-1.c: Adjust testcase. * c-c++-common/goacc/readonly-2.c: New testcase. * gfortran.dg/goacc/readonly-1.f90: Adjust testcase. diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc index 7cf411155c6..42591e4029a 100644 --- a/gcc/c/c-typeck.cc +++ b/gcc/c/c-typeck.cc @@ -14258,6 +14258,8 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH); else OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER); + if (OMP_CLAUSE_MAP_READONLY (c)) + OMP_CLAUSE_MAP_POINTS_TO_READONLY (c2) = 1; OMP_CLAUSE_MAP_IMPLICIT (c2) = OMP_CLAUSE_MAP_IMPLICIT (c); if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER && !c_mark_addressable (t)) diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc index 8fb47fd179e..6ab467e1140 100644 --- a/gcc/cp/semantics.cc +++ b/gcc/cp/semantics.cc @@ -5872,6 +5872,8 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) } else OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER); + if (OMP_CLAUSE_MAP_READONLY (c)) + OMP_CLAUSE_MAP_POINTS_TO_READONLY (c2) = 1; OMP_CLAUSE_MAP_IMPLICIT (c2) = OMP_CLAUSE_MAP_IMPLICIT (c); if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER && !cxx_mark_addressable (t)) diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc index 2253d559f9c..d7cd65af1bb 100644 --- a/gcc/fortran/trans-openmp.cc +++ b/gcc/fortran/trans-openmp.cc @@ -2524,6 +2524,8 @@ gfc_trans_omp_array_section (stmtblock_t *block, gfc_exec_op op, node3 = build_omp_clause (input_location, OMP_CLAUSE_MAP); OMP_CLAUSE_SET_MAP_KIND (node3, ptr_kind); OMP_CLAUSE_DECL (node3) = gfc_conv_descriptor_data_get (decl); + if (n->u.readonly) + OMP_CLAUSE_MAP_POINTS_TO_READONLY (node3) = 1; /* This purposely does not include GOMP_MAP_ALWAYS_POINTER. The extra cast prevents gimplify.cc from recognising it as being part of the struct - and adding an 'alloc: for the 'desc.data' pointer, which @@ -2559,6 +2561,8 @@ gfc_trans_omp_array_section (stmtblock_t *block, gfc_exec_op op, OMP_CLAUSE_MAP); OMP_CLAUSE_SET_MAP_KIND (node3, ptr_kind); OMP_CLAUSE_DECL (node3) = decl; + if (n->u.readonly) + OMP_CLAUSE_MAP_POINTS_TO_READONLY (node3) = 1; } ptr2 = fold_convert (ptrdiff_type_node, ptr2); OMP_CLAUSE_SIZE (node3) = fold_build2 (MINUS_EXPR, ptrdiff_type_node, diff --git a/gcc/gimple-expr.cc b/gcc/gimple-expr.cc index f15cc0ba715..42c0f6469b1 100644 --- a/gcc/gimple-expr.cc +++ b/gcc/gimple-expr.cc @@ -376,6 +376,8 @@ copy_var_decl (tree var, tree name, tree type) DECL_CONTEXT (copy) = DECL_CONTEXT (var); TREE_USED (copy) = 1; DECL_SEEN_IN_BIND_EXPR_P (copy) = 1; + if (VAR_P (var)) + DECL_POINTS_TO_READONLY (copy) = DECL_POINTS_TO_READONLY (var); DECL_ATTRIBUTES (copy) = DECL_ATTRIBUTES (var); if (DECL_USER_ALIGN (var)) { diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index 36e5df050b9..394e40fead2 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -221,6 +221,7 @@ struct gimplify_omp_ctx splay_tree variables; hash_set *privatized_types; tree clauses; + hash_set *pt_readonly_ptrs; /* Iteration variables in an OMP_FOR. */ vec loop_iter_var; location_t location; @@ -628,6 +629,15 @@ internal_get_tmp_var (tree val, gimple_seq *pre_p, gimple_seq *post_p, gimplify_expr (&val, pre_p, post_p, is_gimple_reg_rhs_or_call, fb_rvalue); + bool pt_readonly = false; + if (gimplify_omp_ctxp && gimplify_omp_ctxp->pt_readonly_ptrs) + { + tree ptr = val; + if (TREE_CODE (ptr) == POINTER_PLUS_EXPR) + ptr = TREE_OPERAND (ptr, 0); + pt_readonly = gimplify_omp_ctxp->pt_readonly_ptrs->contains (ptr); + } + if (allow_ssa && gimplify_ctxp->into_ssa && is_gimple_reg_type (TREE_TYPE (val))) @@ -639,9 +649,18 @@ internal_get_tmp_var (tree val, gimple_seq *pre_p, gimple_seq *post_p, if (name) SET_SSA_NAME_VAR_OR_IDENTIFIER (t, create_tmp_var_name (name)); } + if (pt_readonly) + SSA_NAME_POINTS_TO_READONLY_MEMORY (t) = 1; } else - t = lookup_tmp_var (val, is_formal, not_gimple_reg); + { + t = lookup_tmp_var (val, is_formal, not_gimple_reg); + if (pt_readonly) + { + DECL_POINTS_TO_READONLY (t) = 1; + gimplify_omp_ctxp->pt_readonly_ptrs->add (t); + } + } mod = build2 (INIT_EXPR, TREE_TYPE (t), t, unshare_expr (val)); @@ -8906,6 +8925,8 @@ build_omp_struct_comp_nodes (enum tree_code code, tree grp_start, tree grp_end, OMP_CLAUSE_SET_MAP_KIND (c2, mkind); OMP_CLAUSE_DECL (c2) = unshare_expr (OMP_CLAUSE_DECL (grp_end)); OMP_CLAUSE_CHAIN (c2) = NULL_TREE; + if (OMP_CLAUSE_MAP_POINTS_TO_READONLY (grp_end)) + OMP_CLAUSE_MAP_POINTS_TO_READONLY (c2) = 1; tree grp_mid = NULL_TREE; if (OMP_CLAUSE_CHAIN (grp_start) != grp_end) grp_mid = OMP_CLAUSE_CHAIN (grp_start); @@ -11741,6 +11762,16 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, gimplify_omp_ctxp = outer_ctx; } + else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && (code == OACC_PARALLEL + || code == OACC_KERNELS + || code == OACC_SERIAL) + && OMP_CLAUSE_MAP_POINTS_TO_READONLY (c)) + { + if (ctx->pt_readonly_ptrs == NULL) + ctx->pt_readonly_ptrs = new hash_set (); + ctx->pt_readonly_ptrs->add (OMP_CLAUSE_DECL (c)); + } if (notice_outer) goto do_notice; break; diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index b882df048ef..204fc72ca2d 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -14098,6 +14098,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) if (ref_to_array) x = fold_convert_loc (clause_loc, TREE_TYPE (new_var), x); gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue); + if (OMP_CLAUSE_MAP_POINTS_TO_READONLY (c) && VAR_P (x)) + DECL_POINTS_TO_READONLY (x) = 1; if ((is_ref && !ref_to_array) || ref_to_ptr) { diff --git a/gcc/testsuite/c-c++-common/goacc/readonly-1.c b/gcc/testsuite/c-c++-common/goacc/readonly-1.c index 171f96c08db..1f10fd25e46 100644 --- a/gcc/testsuite/c-c++-common/goacc/readonly-1.c +++ b/gcc/testsuite/c-c++-common/goacc/readonly-1.c @@ -19,8 +19,8 @@ int main (void) 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 parallel map\\(readonly,to:\\*s.ptr \\\[len: 64\\\]\\) map\\(pt_readonly,attach_detach:s.ptr \\\[bias: 0\\\]\\) map\\(readonly,to:x\\\[0\\\] \\\[len: 128\\\]\\) map\\(pt_readonly,firstprivate:x \\\[pointer assign, bias: 0\\\]\\)" 1 "original" { target { c } } } } */ +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:\\*NON_LVALUE_EXPR \\\[len: 64\\\]\\) map\\(pt_readonly,attach_detach:s.ptr \\\[bias: 0\\\]\\) map\\(readonly,to:x\\\[0\\\] \\\[len: 128\\\]\\) map\\(pt_readonly,firstprivate:x \\\[pointer assign, bias: 0\\\]\\)" 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/c-c++-common/goacc/readonly-2.c b/gcc/testsuite/c-c++-common/goacc/readonly-2.c new file mode 100644 index 00000000000..d32d3362000 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/readonly-2.c @@ -0,0 +1,15 @@ +/* { dg-additional-options "-O -fdump-tree-fre" } */ + +#pragma acc routine +extern void foo (int *ptr, int val); + +int main (void) +{ + int r, a[32]; + #pragma acc parallel copyin(readonly: a[:32]) copyout(r) + { + foo (a, a[8]); + r = a[8]; + } +} +/* { dg-final { scan-tree-dump-times "r\.\[_0-9\]+ = MEM\\\[\[^_\]+_\[0-9\]+\\(ptro\\)\\\]\\\[8\\\];" 1 "fre1" } } */ diff --git a/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90 b/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90 index 069fec0a0d5..1e5e60f9744 100644 --- a/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90 @@ -20,8 +20,8 @@ program main !$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 parallel map\\(readonly,to:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\) map\\(pt_readonly,alloc:a.0 \\\[pointer assign, bias: \\(integer\\(kind=8\\)\\) parm.*data - \\(integer\\(kind=8\\)\\) a.0\\\]\\) map\\(readonly,to:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\) map\\(pt_readonly,alloc:b \\\[pointer assign, bias: \\(integer\\(kind=8\\)\\) parm.*data - \\(integer\\(kind=8\\)\\) b\\\]\\)" 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\\(pt_readonly,alloc:a \\\[pointer assign, bias: \\(integer\\(kind=8\\)\\) parm.*data - \\(integer\\(kind=8\\)\\) &a\\\]\\) map\\(readonly,to:b\\\[\\(\\(integer\\(kind=8\\)\\) parm.*data - \\(integer\\(kind=8\\)\\) &b\\) / 4\\\] \\\[len: .+\\\]\\) map\\(pt_readonly,alloc:b \\\[pointer assign, bias: \\(integer\\(kind=8\\)\\) parm.*data - \\(integer\\(kind=8\\)\\) &b\\\]\\)" 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 9604c3eecc5..1a8b121f30b 100644 --- a/gcc/tree-pretty-print.cc +++ b/gcc/tree-pretty-print.cc @@ -907,6 +907,8 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) pp_string (pp, "map("); if (OMP_CLAUSE_MAP_READONLY (clause)) pp_string (pp, "readonly,"); + if (OMP_CLAUSE_MAP_POINTS_TO_READONLY (clause)) + pp_string (pp, "pt_readonly,"); switch (OMP_CLAUSE_MAP_KIND (clause)) { case GOMP_MAP_ALLOC: @@ -3436,6 +3438,8 @@ dump_generic_node (pretty_printer *pp, tree node, int spc, dump_flags_t flags, pp_string (pp, "(D)"); if (SSA_NAME_OCCURS_IN_ABNORMAL_PHI (node)) pp_string (pp, "(ab)"); + if (SSA_NAME_POINTS_TO_READONLY_MEMORY (node)) + pp_string (pp, "(ptro)"); break; case WITH_SIZE_EXPR: diff --git a/gcc/tree-ssanames.cc b/gcc/tree-ssanames.cc index 23387b90fe3..32d35a29dfc 100644 --- a/gcc/tree-ssanames.cc +++ b/gcc/tree-ssanames.cc @@ -402,6 +402,9 @@ make_ssa_name_fn (struct function *fn, tree var, gimple *stmt, else SSA_NAME_RANGE_INFO (t) = NULL; + if (VAR_P (var) && DECL_POINTS_TO_READONLY (var)) + SSA_NAME_POINTS_TO_READONLY_MEMORY (t) = 1; + SSA_NAME_IN_FREE_LIST (t) = 0; SSA_NAME_IS_DEFAULT_DEF (t) = 0; init_ssa_name_imm_use (t); diff --git a/gcc/tree.h b/gcc/tree.h index ac563de1fc3..880ffb367a3 100644 --- a/gcc/tree.h +++ b/gcc/tree.h @@ -1021,6 +1021,13 @@ extern void omp_clause_range_check_failed (const_tree, const char *, int, #define DECL_HIDDEN_STRING_LENGTH(NODE) \ (TREE_CHECK (NODE, PARM_DECL)->decl_common.decl_nonshareable_flag) +/* In a VAR_DECL, set for variables regarded as pointing to memory not written + to. SSA_NAME_POINTS_TO_READONLY_MEMORY gets set for SSA_NAMEs created from + such VAR_DECLs. Currently used by OpenACC 'readonly' modifier in copyin + clauses. */ +#define DECL_POINTS_TO_READONLY(NODE) \ + (TREE_CHECK (NODE, VAR_DECL)->decl_common.decl_not_flexarray) + /* In a CALL_EXPR, means that the call is the jump from a thunk to the thunked-to function. Be careful to avoid using this macro when one of the next two applies instead. */ @@ -1815,6 +1822,10 @@ class auto_suppress_location_wrappers #define OMP_CLAUSE_MAP_READONLY(NODE) \ TREE_READONLY (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)) +/* Set if 'OMP_CLAUSE_DECL (NODE)' points to read-only memory. */ +#define OMP_CLAUSE_MAP_POINTS_TO_READONLY(NODE) \ + TREE_CONSTANT (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_))