From patchwork Tue May 11 08:56:31 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1476980 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=8.43.85.97; helo=sourceware.org; envelope-from=gcc-patches-bounces@gcc.gnu.org; receiver=) Received: from sourceware.org (ip-8-43-85-97.sourceware.org [8.43.85.97]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature RSA-PSS (4096 bits) server-digest SHA256) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 4FfWyJ1H5xz9sjJ for ; Tue, 11 May 2021 18:57:12 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 55A6A3886C55; Tue, 11 May 2021 08:57:02 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa2.mentor.iphmx.com (esa2.mentor.iphmx.com [68.232.141.98]) by sourceware.org (Postfix) with ESMTPS id 1737F3857801; Tue, 11 May 2021 08:56:59 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org 1737F3857801 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=Julian_Brown@mentor.com IronPort-SDR: 81zAQP3hCe89jzFacizRL69QEse/R1hs36sfZLx4y+RcqUDzPGOPj6niPPIkM8MHGZKn+550Sl i8Ze+dabOQmlTU8fhBlbrI4mX6EGDGGeiyp3MB9v6GI37GrSsR6opBdB5MRKtJmcnXlDQrYH/y HhHSleAKtGZdGlHVMRQYOT1rIDg6H6MYsagv+wEY2BmcaXuNN6XJXPM9bGPnB4s78iqhsE/u7l ejoCsOIP4Uww7j6DidQVTlibVBjVlImv7WCkcJmHaYZ36Zdn3Ecse4XNxV5H3kXuykxVedUG7e zLI= X-IronPort-AV: E=Sophos;i="5.82,290,1613462400"; d="scan'208";a="61053240" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa2.mentor.iphmx.com with ESMTP; 11 May 2021 00:56:59 -0800 IronPort-SDR: VFFXf3ROJCF2ENJ15yXmSDvMGYQu1UpHhEdrKOue0waJak9iN/8tcws41V5158Y21qw5m/n8Jj mOhITo+Px9AgFyvmoWV4kM7om7VVPTlklqQmArGkr4kmfXbV3AH3HgvMeqgyXb3y32vIug9waa DZXwrekvc+zZRijt29qy9puEArsHcvi0fNFt3SIaZNYnnN2C1/9GSQAe8qCmy1W76DBct9e2dB wBfVCheRsfFqlGwqFOfsi3XfubpfCG2N2gF3H+hxOeIkb2TgX0MVD129Zv/THtJ0v5cmydsDGk /jY= From: Julian Brown To: Subject: [PATCH 1/7] [og10] Unify ARRAY_REF/INDIRECT_REF stripping code in extract_base_bit_offset Date: Tue, 11 May 2021 01:56:31 -0700 Message-ID: <23c6b39543ca5a94d59682cc66e136109da086ea.1620721888.git.julian@codesourcery.com> X-Mailer: git-send-email 2.29.2 In-Reply-To: References: MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-09.mgc.mentorg.com (139.181.222.9) To SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) X-Spam-Status: No, score=-12.0 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, KAM_SHORT, RCVD_IN_DNSWL_NONE, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.2 X-Spam-Checker-Version: SpamAssassin 3.4.2 (2018-09-13) 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: , Cc: Jakub Jelinek , Tobias Burnus , fortran@gcc.gnu.org, Thomas Schwinge Errors-To: gcc-patches-bounces@gcc.gnu.org Sender: "Gcc-patches" For historical reasons, it seems that extract_base_bit_offset unnecessarily used two different ways to strip ARRAY_REF/INDIRECT_REF nodes from component accesses. I verified that the two ways of performing the operation gave the same results across the whole testsuite (and several additional benchmarks). The code was like this since an earlier "mechanical" refactoring by me, first posted here: https://gcc.gnu.org/pipermail/gcc-patches/2018-November/510503.html It was never clear to me if there was an important semantic difference between the two ways of stripping the base before calling get_inner_reference, but it appears that there is not, so one can go away. 2021-05-11 Julian Brown gcc/ * gimplify.c (extract_base_bit_offset): Unify ARRAY_REF/INDIRECT_REF stripping code in first call/subsequent call cases. --- gcc/gimplify.c | 32 +++++++++++--------------------- 1 file changed, 11 insertions(+), 21 deletions(-) diff --git a/gcc/gimplify.c b/gcc/gimplify.c index ba071e8ae68..b36b961bf26 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -8340,31 +8340,21 @@ extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp, poly_offset_int poffset; if (base_ref) - { - *base_ref = NULL_TREE; - - while (TREE_CODE (base) == ARRAY_REF) - base = TREE_OPERAND (base, 0); + *base_ref = NULL_TREE; - if (TREE_CODE (base) == INDIRECT_REF) - base = TREE_OPERAND (base, 0); - } - else + if (TREE_CODE (base) == ARRAY_REF) { - if (TREE_CODE (base) == ARRAY_REF) - { - while (TREE_CODE (base) == ARRAY_REF) - base = TREE_OPERAND (base, 0); - if (TREE_CODE (base) != COMPONENT_REF - || TREE_CODE (TREE_TYPE (base)) != ARRAY_TYPE) - return NULL_TREE; - } - else if (TREE_CODE (base) == INDIRECT_REF - && TREE_CODE (TREE_OPERAND (base, 0)) == COMPONENT_REF - && (TREE_CODE (TREE_TYPE (TREE_OPERAND (base, 0))) - == REFERENCE_TYPE)) + while (TREE_CODE (base) == ARRAY_REF) base = TREE_OPERAND (base, 0); + if (TREE_CODE (base) != COMPONENT_REF + || TREE_CODE (TREE_TYPE (base)) != ARRAY_TYPE) + return NULL_TREE; } + else if (TREE_CODE (base) == INDIRECT_REF + && TREE_CODE (TREE_OPERAND (base, 0)) == COMPONENT_REF + && (TREE_CODE (TREE_TYPE (TREE_OPERAND (base, 0))) + == REFERENCE_TYPE)) + base = TREE_OPERAND (base, 0); base = get_inner_reference (base, &bitsize, &bitpos, &offset, &mode, &unsignedp, &reversep, &volatilep); From patchwork Tue May 11 08:56:32 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1476981 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=8.43.85.97; helo=sourceware.org; envelope-from=gcc-patches-bounces@gcc.gnu.org; receiver=) Received: from sourceware.org (ip-8-43-85-97.sourceware.org [8.43.85.97]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature RSA-PSS (4096 bits) server-digest SHA256) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 4FfWyR5bq5z9sT6 for ; Tue, 11 May 2021 18:57:19 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 0EAE13886C5F; Tue, 11 May 2021 08:57:09 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa2.mentor.iphmx.com (esa2.mentor.iphmx.com [68.232.141.98]) by sourceware.org (Postfix) with ESMTPS id 573F03839C4D; Tue, 11 May 2021 08:57:04 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org 573F03839C4D Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=Julian_Brown@mentor.com IronPort-SDR: FVHUaK4b0CbYSP7neaCAf9ydJm6+hBPp+d5CELNYHqnu1Whq42lNAXl63yxeO6ahrxn3ooGb7N L+olIFHGD/dXa/nAqsGb+lZWoHXp9g3RvqhMpBBaVF+g0GAjo2BKs3eQ3roJUescgQtxZy899D rm0jKhvJMxzhUFCjQzlcK1Egz8lD+3BB5hASREcZQ5roFmlhTNNCFziOD+xWKa8ZIrZT1C+SGb z3/iTMHZcw0fKGJz3hYxCQ/E50ITBlKw8Qg0NCIpfEh6ka/jQM9scb+zrf4qBN/4vCCpP9C3NB Hjg= X-IronPort-AV: E=Sophos;i="5.82,290,1613462400"; d="scan'208";a="61053249" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa2.mentor.iphmx.com with ESMTP; 11 May 2021 00:57:03 -0800 IronPort-SDR: QRFikACzWZ4VZzNWVrA1Itg5tV7pkErwR7Pe9A+Of/R5PouNugfrQsC26o0zmDIMuNYpBhYp7s ZH5cqD7nk83C08kx24KjFw7rwDY/iM80Vonrp/0nTeQLTUyHj1SDzG7iMP4jGs5aInAaXiVXn/ siK5FVxll95WoXbepJcreScfb2WHhDiyk1IYS3N1fEXoYtt2nqHvBTpRw2zXysQCcgcmRUPK5D omueMMFfIj9t6Jq/9YcqkricsY5AdWH8x+dEfrk0S9BchIWWfKujeG6eVohHTDY/dCaLbUKlJh k7Q= From: Julian Brown To: Subject: [PATCH 2/7] [og10] Refactor struct lowering for OpenMP/OpenACC in gimplify.c Date: Tue, 11 May 2021 01:56:32 -0700 Message-ID: <0102d1f2aa246242001e6a9fa0fe371beb95108f.1620721888.git.julian@codesourcery.com> X-Mailer: git-send-email 2.29.2 In-Reply-To: References: MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-09.mgc.mentorg.com (139.181.222.9) To SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) X-Spam-Status: No, score=-11.6 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_ASCII_DIVIDERS, KAM_DMARC_STATUS, KAM_SHORT, RCVD_IN_DNSWL_NONE, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.2 X-Spam-Checker-Version: SpamAssassin 3.4.2 (2018-09-13) 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: , Cc: Jakub Jelinek , Tobias Burnus , fortran@gcc.gnu.org, Thomas Schwinge Errors-To: gcc-patches-bounces@gcc.gnu.org Sender: "Gcc-patches" This patch is a second attempt at refactoring struct component mapping handling for OpenACC/OpenMP during gimplification, after the patch I posted here: https://gcc.gnu.org/pipermail/gcc-patches/2018-November/510503.html And improved here, post-review: https://gcc.gnu.org/pipermail/gcc-patches/2019-November/533394.html This patch goes further, in that the struct-handling code is outlined into its own function (to create the "GOMP_MAP_STRUCT" node and the sorted list of nodes immediately following it, from a set of mappings of components of a given struct or derived type). I've also gone through the list-handling code and attempted to add comments documenting how it works to the best of my understanding, and broken out a couple of helper functions in order to (hopefully) have the code self-document better also. 2021-05-11 Julian Brown gcc/ * gimplify.c (insert_struct_comp_map): Refactor function into... (build_struct_comp_nodes): This new function. Remove list handling and improve self-documentation. (insert_node_after, move_node_after, move_nodes_after, move_concat_nodes_after): New helper functions. (build_struct_group): New function to build up GOMP_MAP_STRUCT node groups to map struct components. Outlined from... (gimplify_scan_omp_clauses): Here. --- gcc/gimplify.c | 846 +++++++++++++++++++++++++++++++------------------ 1 file changed, 540 insertions(+), 306 deletions(-) diff --git a/gcc/gimplify.c b/gcc/gimplify.c index b36b961bf26..ad192b27208 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -8253,73 +8253,66 @@ gimplify_omp_depend (tree *list_p, gimple_seq *pre_p) return 1; } -/* Insert a GOMP_MAP_ALLOC or GOMP_MAP_RELEASE node following a - GOMP_MAP_STRUCT mapping. C is an always_pointer mapping. STRUCT_NODE is - the struct node to insert the new mapping after (when the struct node is - initially created). PREV_NODE is the first of two or three mappings for a - pointer, and is either: - - the node before C, when a pair of mappings is used, e.g. for a C/C++ - array section. - - not the node before C. This is true when we have a reference-to-pointer - type (with a mapping for the reference and for the pointer), or for - Fortran derived-type mappings with a GOMP_MAP_TO_PSET. - If SCP is non-null, the new node is inserted before *SCP. - if SCP is null, the new node is inserted before PREV_NODE. - The return type is: - - PREV_NODE, if SCP is non-null. - - The newly-created ALLOC or RELEASE node, if SCP is null. - - The second newly-created ALLOC or RELEASE node, if we are mapping a - reference to a pointer. */ +/* For a set of mappings describing an array section pointed to by a struct + (or derived type, etc.) component, create an "alloc" or "release" node to + insert into a list following a GOMP_MAP_STRUCT node. For some types of + mapping (e.g. Fortran arrays with descriptors), an additional mapping may + be created that is inserted into the list of mapping nodes attached to the + directive being processed -- not part of the sorted list of nodes after + GOMP_MAP_STRUCT. + + CODE is the code of the directive being processed. GRP_START and GRP_END + are the first and last of two or three nodes representing this array section + mapping (e.g. a data movement node like GOMP_MAP_{TO,FROM}, optionally a + GOMP_MAP_TO_PSET, and finally a GOMP_MAP_ALWAYS_POINTER). EXTRA_NODE is + filled with the additional node described above, if needed. + + This function does not add the new nodes to any lists itself. It is the + responsibility of the caller to do that. */ static tree -insert_struct_comp_map (enum tree_code code, tree c, tree struct_node, - tree prev_node, tree *scp) +build_struct_comp_nodes (enum tree_code code, tree grp_start, tree grp_end, + tree *extra_node) { enum gomp_map_kind mkind = (code == OMP_TARGET_EXIT_DATA || code == OACC_EXIT_DATA) ? GOMP_MAP_RELEASE : GOMP_MAP_ALLOC; - tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); - tree cl = scp ? prev_node : c2; + gcc_assert (grp_start != grp_end); + + tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (grp_end), OMP_CLAUSE_MAP); OMP_CLAUSE_SET_MAP_KIND (c2, mkind); - OMP_CLAUSE_DECL (c2) = unshare_expr (OMP_CLAUSE_DECL (c)); - OMP_CLAUSE_CHAIN (c2) = scp ? *scp : prev_node; - if (OMP_CLAUSE_CHAIN (prev_node) != c - && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (prev_node)) == OMP_CLAUSE_MAP - && (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (prev_node)) - == GOMP_MAP_TO_PSET)) - OMP_CLAUSE_SIZE (c2) = OMP_CLAUSE_SIZE (OMP_CLAUSE_CHAIN (prev_node)); + OMP_CLAUSE_DECL (c2) = unshare_expr (OMP_CLAUSE_DECL (grp_end)); + OMP_CLAUSE_CHAIN (c2) = NULL_TREE; + tree grp_mid = NULL_TREE; + if (OMP_CLAUSE_CHAIN (grp_start) != grp_end) + grp_mid = OMP_CLAUSE_CHAIN (grp_start); + + if (grp_mid + && OMP_CLAUSE_CODE (grp_mid) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_KIND (grp_mid) == GOMP_MAP_TO_PSET) + OMP_CLAUSE_SIZE (c2) = OMP_CLAUSE_SIZE (grp_mid); else OMP_CLAUSE_SIZE (c2) = TYPE_SIZE_UNIT (ptr_type_node); - if (struct_node) - OMP_CLAUSE_CHAIN (struct_node) = c2; - - /* We might need to create an additional mapping if we have a reference to a - pointer (in C++). Don't do this if we have something other than a - GOMP_MAP_ALWAYS_POINTER though, i.e. a GOMP_MAP_TO_PSET. */ - if (OMP_CLAUSE_CHAIN (prev_node) != c - && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (prev_node)) == OMP_CLAUSE_MAP - && ((OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (prev_node)) - == GOMP_MAP_ALWAYS_POINTER) - || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (prev_node)) - == GOMP_MAP_ATTACH_DETACH))) - { - tree c4 = OMP_CLAUSE_CHAIN (prev_node); - tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); + + if (grp_mid + && OMP_CLAUSE_CODE (grp_mid) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (grp_mid) == GOMP_MAP_ALWAYS_POINTER + || OMP_CLAUSE_MAP_KIND (grp_mid) == GOMP_MAP_ATTACH_DETACH)) + { + tree c3 + = build_omp_clause (OMP_CLAUSE_LOCATION (grp_end), OMP_CLAUSE_MAP); OMP_CLAUSE_SET_MAP_KIND (c3, mkind); - OMP_CLAUSE_DECL (c3) = unshare_expr (OMP_CLAUSE_DECL (c4)); + OMP_CLAUSE_DECL (c3) = unshare_expr (OMP_CLAUSE_DECL (grp_mid)); OMP_CLAUSE_SIZE (c3) = TYPE_SIZE_UNIT (ptr_type_node); - OMP_CLAUSE_CHAIN (c3) = prev_node; - if (!scp) - OMP_CLAUSE_CHAIN (c2) = c3; - else - cl = c3; - } + OMP_CLAUSE_CHAIN (c3) = NULL_TREE; - if (scp) - *scp = c2; + *extra_node = c3; + } + else + *extra_node = NULL_TREE; - return cl; + return c2; } /* Strip ARRAY_REFS or an indirect ref off BASE, find the containing object, @@ -8495,6 +8488,486 @@ omp_target_reorder_clauses (tree *list_p) } } +/* Link node NEWNODE so it is pointed to by chain INSERT_AT. NEWNODE's chain + is linked to the previous node pointed to by INSERT_AT. */ + +static tree * +insert_node_after (tree newnode, tree *insert_at) +{ + OMP_CLAUSE_CHAIN (newnode) = *insert_at; + *insert_at = newnode; + return &OMP_CLAUSE_CHAIN (newnode); +} + +/* Move NODE (which is currently pointed to by the chain OLD_POS) so it is + pointed to by chain MOVE_AFTER instead. */ + +static void +move_node_after (tree node, tree *old_pos, tree *move_after) +{ + gcc_assert (node == *old_pos); + *old_pos = OMP_CLAUSE_CHAIN (node); + OMP_CLAUSE_CHAIN (node) = *move_after; + *move_after = node; +} + +/* Move nodes from FIRST_PTR (pointed to by previous node's chain) to + LAST_NODE to after MOVE_AFTER chain. Similar to below function, but no + new nodes are prepended to the list before splicing into the new position. + Return the position we should continue scanning the list at, or NULL to + stay where we were. */ + +static tree * +move_nodes_after (tree *first_ptr, tree last_node, tree *move_after) +{ + if (first_ptr == move_after) + return NULL; + + tree tmp = *first_ptr; + *first_ptr = OMP_CLAUSE_CHAIN (last_node); + OMP_CLAUSE_CHAIN (last_node) = *move_after; + *move_after = tmp; + + return first_ptr; +} + +/* Concatenate two lists described by [FIRST_NEW, LAST_NEW_TAIL] and + [FIRST_PTR, LAST_NODE], and insert them in the OMP clause list after chain + pointer MOVE_AFTER. + + The latter list was previously part of the OMP clause list, and the former + (prepended) part is comprised of new nodes. + + We start with a list of nodes starting with a struct mapping node. We + rearrange the list so that new nodes starting from FIRST_NEW and whose last + node's chain is LAST_NEW_TAIL comes directly after MOVE_AFTER, followed by + the group of mapping nodes we are currently processing (from the chain + FIRST_PTR to LAST_NODE). The return value is the pointer to the next chain + we should continue processing from, or NULL to stay where we were. + + The transformation (in the case where MOVE_AFTER and FIRST_PTR are + different) is worked through below. Here we are processing LAST_NODE, and + FIRST_PTR points at the preceding mapping clause: + + #. mapping node chain + --------------------------------------------------- + A. struct_node [->B] + B. comp_1 [->C] + C. comp_2 [->D (move_after)] + D. map_to_3 [->E] + E. attach_3 [->F (first_ptr)] + F. map_to_4 [->G (continue_at)] + G. attach_4 (last_node) [->H] + H. ... + + *last_new_tail = *first_ptr; + + I. new_node (first_new) [->F (last_new_tail)] + + *first_ptr = OMP_CLAUSE_CHAIN (last_node) + + #. mapping node chain + ---------------------------------------------------- + A. struct_node [->B] + B. comp_1 [->C] + C. comp_2 [->D (move_after)] + D. map_to_3 [->E] + E. attach_3 [->H (first_ptr)] + F. map_to_4 [->G (continue_at)] + G. attach_4 (last_node) [->H] + H. ... + + I. new_node (first_new) [->F (last_new_tail)] + + OMP_CLAUSE_CHAIN (last_node) = *move_after; + + #. mapping node chain + --------------------------------------------------- + A. struct_node [->B] + B. comp_1 [->C] + C. comp_2 [->D (move_after)] + D. map_to_3 [->E] + E. attach_3 [->H (continue_at)] + F. map_to_4 [->G] + G. attach_4 (last_node) [->D] + H. ... + + I. new_node (first_new) [->F (last_new_tail)] + + *move_after = first_new; + + #. mapping node chain + --------------------------------------------------- + A. struct_node [->B] + B. comp_1 [->C] + C. comp_2 [->I (move_after)] + D. map_to_3 [->E] + E. attach_3 [->H (continue_at)] + F. map_to_4 [->G] + G. attach_4 (last_node) [->D] + H. ... + I. new_node (first_new) [->F (last_new_tail)] + + or, in order: + + #. mapping node chain + --------------------------------------------------- + A. struct_node [->B] + B. comp_1 [->C] + C. comp_2 [->I (move_after)] + I. new_node (first_new) [->F (last_new_tail)] + F. map_to_4 [->G] + G. attach_4 (last_node) [->D] + D. map_to_3 [->E] + E. attach_3 [->H (continue_at)] + H. ... +*/ + +static tree * +move_concat_nodes_after (tree first_new, tree *last_new_tail, tree *first_ptr, + tree last_node, tree *move_after) +{ + tree *continue_at = NULL; + *last_new_tail = *first_ptr; + if (first_ptr == move_after) + *move_after = first_new; + else + { + *first_ptr = OMP_CLAUSE_CHAIN (last_node); + continue_at = first_ptr; + OMP_CLAUSE_CHAIN (last_node) = *move_after; + *move_after = first_new; + } + return continue_at; +} + +/* Mapping struct members causes an additional set of nodes to be created, + starting with GOMP_MAP_STRUCT followed by a number of mappings equal to the + number of members being mapped, in order of ascending position (address or + bitwise). + + We scan through the list of mapping clauses, calling this function for each + struct member mapping we find, and build up the list of mappings after the + initial GOMP_MAP_STRUCT node. For pointer members, these will be + newly-created ALLOC nodes. For non-pointer members, the existing mapping is + moved into place in the sorted list. + + struct { + int *a; + int *b; + int c; + int *d; + }; + + #pragma (acc|omp directive) copy(struct.a[0:n], struct.b[0:n], struct.c, + struct.d[0:n]) + + GOMP_MAP_STRUCT (4) + [GOMP_MAP_FIRSTPRIVATE_REFERENCE -- for refs to structs] + GOMP_MAP_ALLOC (struct.a) + GOMP_MAP_ALLOC (struct.b) + GOMP_MAP_TO (struct.c) + GOMP_MAP_ALLOC (struct.d) + ... + + In the case where we are mapping references to pointers, or in Fortran if + we are mapping an array with a descriptor, additional nodes may be created + after the struct node list also. + + The return code is: + - DECL, if we just created the initial GOMP_MAP_STRUCT node. + - NULL_TREE, if we inserted the new struct member successfully. + - error_mark_node if an error occurred. + + *CONT is set to TRUE if we should skip further processing and move to the + next node. PREV_LIST_P and LIST_P may be modified by the function when a + list rearrangement has taken place. */ + +static tree +build_struct_group (struct gimplify_omp_ctx *ctx, + enum omp_region_type region_type, enum tree_code code, + tree decl, tree *pd, bool component_ref_p, + unsigned int *flags, tree c, + hash_map *&struct_map_to_clause, + tree *&prev_list_p, tree *&list_p, gimple_seq *pre_p, + bool *cont) +{ + poly_offset_int coffset; + poly_int64 cbitpos; + tree base_ref; + + tree base = extract_base_bit_offset (OMP_CLAUSE_DECL (c), &base_ref, + &cbitpos, &coffset); + + gcc_assert (base == decl); + + splay_tree_node n + = (DECL_P (decl) + ? splay_tree_lookup (ctx->variables, (splay_tree_key) decl) + : NULL); + bool ptr = (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER); + bool attach_detach = (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH); + bool attach = (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH); + bool has_attachments = false; + /* For OpenACC, pointers in structs should trigger an attach action. */ + if (attach_detach + && ((region_type & (ORT_ACC | ORT_TARGET | ORT_TARGET_DATA)) + || code == OMP_TARGET_ENTER_DATA + || code == OMP_TARGET_EXIT_DATA)) + { + /* Turn a GOMP_MAP_ATTACH_DETACH clause into a GOMP_MAP_ATTACH or + GOMP_MAP_DETACH clause after we have detected a case that needs a + GOMP_MAP_STRUCT mapping added. */ + gomp_map_kind k + = ((code == OACC_EXIT_DATA || code == OMP_TARGET_EXIT_DATA) + ? GOMP_MAP_DETACH : GOMP_MAP_ATTACH); + OMP_CLAUSE_SET_MAP_KIND (c, k); + has_attachments = true; + } + if ((DECL_P (decl) && (n == NULL || (n->value & GOVD_MAP) == 0)) + || (!DECL_P (decl) + && (!struct_map_to_clause + || struct_map_to_clause->get (decl) == NULL))) + { + tree l = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); + gomp_map_kind k = attach ? GOMP_MAP_FORCE_PRESENT : GOMP_MAP_STRUCT; + + OMP_CLAUSE_SET_MAP_KIND (l, k); + if (base_ref) + OMP_CLAUSE_DECL (l) = unshare_expr (base_ref); + else + { + OMP_CLAUSE_DECL (l) = unshare_expr (decl); + if (!DECL_P (OMP_CLAUSE_DECL (l)) + && (gimplify_expr (&OMP_CLAUSE_DECL (l), pre_p, NULL, + is_gimple_lvalue, fb_lvalue) == GS_ERROR)) + return error_mark_node; + } + OMP_CLAUSE_SIZE (l) + = (!attach ? size_int (1) + : (DECL_P (OMP_CLAUSE_DECL (l)) + ? DECL_SIZE_UNIT (OMP_CLAUSE_DECL (l)) + : TYPE_SIZE_UNIT (TREE_TYPE (OMP_CLAUSE_DECL (l))))); + if (struct_map_to_clause == NULL) + struct_map_to_clause = new hash_map; + struct_map_to_clause->put (decl, l); + if (ptr || attach_detach) + { + tree extra_node; + tree alloc_node + = build_struct_comp_nodes (code, *prev_list_p, c, &extra_node); + OMP_CLAUSE_CHAIN (l) = alloc_node; + if (extra_node) + { + OMP_CLAUSE_CHAIN (extra_node) = *prev_list_p; + OMP_CLAUSE_CHAIN (alloc_node) = extra_node; + } + else + OMP_CLAUSE_CHAIN (alloc_node) = *prev_list_p; + *prev_list_p = l; + prev_list_p = NULL; + } + else + list_p = insert_node_after (l, list_p); + if (base_ref && code == OMP_TARGET) + { + tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); + enum gomp_map_kind mkind = GOMP_MAP_FIRSTPRIVATE_REFERENCE; + OMP_CLAUSE_SET_MAP_KIND (c2, mkind); + OMP_CLAUSE_DECL (c2) = decl; + OMP_CLAUSE_SIZE (c2) = size_zero_node; + OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (l); + OMP_CLAUSE_CHAIN (l) = c2; + } + *flags = GOVD_MAP | GOVD_EXPLICIT; + if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) || ptr || attach_detach) + *flags |= GOVD_SEEN; + if (has_attachments) + *flags |= GOVD_MAP_HAS_ATTACHMENTS; + + /* If this is a *pointer-to-struct expression, make sure a + firstprivate map of the base-pointer exists. */ + if (component_ref_p + && ((TREE_CODE (decl) == MEM_REF + && integer_zerop (TREE_OPERAND (decl, 1))) + || INDIRECT_REF_P (decl)) + && DECL_P (TREE_OPERAND (decl, 0)) + && !splay_tree_lookup (ctx->variables, + ((splay_tree_key) TREE_OPERAND (decl, 0)))) + { + decl = TREE_OPERAND (decl, 0); + tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); + enum gomp_map_kind mkind = GOMP_MAP_FIRSTPRIVATE_POINTER; + OMP_CLAUSE_SET_MAP_KIND (c2, mkind); + OMP_CLAUSE_DECL (c2) = decl; + OMP_CLAUSE_SIZE (c2) = size_zero_node; + OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c); + OMP_CLAUSE_CHAIN (c) = c2; + } + + return decl; + } + else if (struct_map_to_clause) + { + tree *osc = struct_map_to_clause->get (decl); + tree *sc = NULL, *scp = NULL; + if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) || ptr || attach_detach) + n->value |= GOVD_SEEN; + sc = &OMP_CLAUSE_CHAIN (*osc); + if (*sc != c + && OMP_CLAUSE_MAP_KIND (*sc) == GOMP_MAP_FIRSTPRIVATE_REFERENCE) + sc = &OMP_CLAUSE_CHAIN (*sc); + for (; *sc != c; sc = &OMP_CLAUSE_CHAIN (*sc)) + { + if ((ptr || attach_detach) && sc == prev_list_p) + break; + else if (TREE_CODE (OMP_CLAUSE_DECL (*sc)) != COMPONENT_REF + && TREE_CODE (OMP_CLAUSE_DECL (*sc)) != INDIRECT_REF + && TREE_CODE (OMP_CLAUSE_DECL (*sc)) != ARRAY_REF) + break; + else + { + tree sc_decl = OMP_CLAUSE_DECL (*sc); + poly_offset_int offset; + poly_int64 bitpos; + tree base = extract_base_bit_offset (sc_decl, NULL, &bitpos, + &offset); + if (base != decl) + break; + if (scp) + continue; + if ((region_type & ORT_ACC) != 0) + { + /* This duplicate checking code is currently only + enabled for OpenACC. */ + tree d1 = OMP_CLAUSE_DECL (*sc); + tree d2 = OMP_CLAUSE_DECL (c); + while (TREE_CODE (d1) == ARRAY_REF) + d1 = TREE_OPERAND (d1, 0); + while (TREE_CODE (d2) == ARRAY_REF) + d2 = TREE_OPERAND (d2, 0); + if (TREE_CODE (d1) == INDIRECT_REF) + d1 = TREE_OPERAND (d1, 0); + if (TREE_CODE (d2) == INDIRECT_REF) + d2 = TREE_OPERAND (d2, 0); + while (TREE_CODE (d1) == COMPONENT_REF) + if (TREE_CODE (d2) == COMPONENT_REF + && TREE_OPERAND (d1, 1) == TREE_OPERAND (d2, 1)) + { + d1 = TREE_OPERAND (d1, 0); + d2 = TREE_OPERAND (d2, 0); + } + else + break; + if (d1 == d2) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qE appears more than once in map clauses", + OMP_CLAUSE_DECL (c)); + return error_mark_node; + } + } + if (maybe_lt (coffset, offset) + || (known_eq (coffset, offset) + && maybe_lt (cbitpos, bitpos))) + { + if (ptr || attach_detach) + scp = sc; + else + break; + } + } + } + + if (!attach) + OMP_CLAUSE_SIZE (*osc) + = size_binop (PLUS_EXPR, OMP_CLAUSE_SIZE (*osc), size_one_node); + if (ptr || attach_detach) + { + tree cl = NULL_TREE, extra_node; + tree alloc_node = build_struct_comp_nodes (code, *prev_list_p, c, + &extra_node); + tree *tail_chain = NULL; + + /* Here, we have: + + c : the currently-processed node. + prev_list_p : pointer to the first node in a pointer mapping group + up to and including C. + sc : pointer to the chain for the end of the struct component + list. + scp : pointer to the chain for the sorted position at which we + should insert in the middle of the struct component list + (else NULL to insert at end). + alloc_node : the "alloc" node for the structure (pointer-type) + component. We insert at SCP (if present), else SC + (the end of the struct component list). + extra_node : a newly-synthesized node for an additional indirect + pointer mapping or a Fortran pointer set, if needed. + cl : first node to prepend before prev_list_p. + tail_chain : pointer to chain of last prepended node. + + The general idea is we move the nodes for this struct mapping + together: the alloc node goes into the sorted list directly after + the struct mapping, and any extra nodes (together with the nodes + mapping arrays pointed to by struct components) get moved after + that list. When SCP is NULL, we insert the nodes at SC, i.e. at + the end of the struct component mapping list. It's important that + the alloc_node comes first in that case because it's part of the + sorted component mapping list (but subsequent nodes are not!). */ + + if (scp) + insert_node_after (alloc_node, scp); + + /* Make [cl,tail_chain] a list of the alloc node (if we haven't + already inserted it) and the extra_node (if it is present). The + list can be empty if we added alloc_node above and there is no + extra node. */ + if (scp && extra_node) + { + cl = extra_node; + tail_chain = &OMP_CLAUSE_CHAIN (extra_node); + } + else if (extra_node) + { + OMP_CLAUSE_CHAIN (alloc_node) = extra_node; + cl = alloc_node; + tail_chain = &OMP_CLAUSE_CHAIN (extra_node); + } + else if (!scp) + { + cl = alloc_node; + tail_chain = &OMP_CLAUSE_CHAIN (alloc_node); + } + + tree *continue_at + = cl ? move_concat_nodes_after (cl, tail_chain, prev_list_p, c, sc) + : move_nodes_after (prev_list_p, c, sc); + + prev_list_p = NULL; + + if (continue_at) + { + list_p = continue_at; + *cont = true; + } + } + else if (*sc != c) + { + if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, fb_lvalue) + == GS_ERROR) + return error_mark_node; + /* In the non-pointer case, the mapping clause itself is moved into + the correct position in the struct component list, which in this + case is just SC. */ + move_node_after (c, list_p, sc); + *cont = true; + } + } + return NULL_TREE; +} + /* Scan the OMP clauses in *LIST_P, installing mappings into a new and previous omp contexts. */ @@ -9235,263 +9708,24 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, } } } - - poly_offset_int offset1; - poly_int64 bitpos1; - tree base_ref; - - tree base - = extract_base_bit_offset (OMP_CLAUSE_DECL (c), &base_ref, - &bitpos1, &offset1); - - gcc_assert (base == decl); - - splay_tree_node n - = (DECL_P (decl) - ? splay_tree_lookup (ctx->variables, - (splay_tree_key) decl) - : NULL); - bool ptr = (OMP_CLAUSE_MAP_KIND (c) - == GOMP_MAP_ALWAYS_POINTER); - bool attach_detach = (OMP_CLAUSE_MAP_KIND (c) - == GOMP_MAP_ATTACH_DETACH); - bool attach = OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH - || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH; - bool has_attachments = false; - /* For OpenACC, pointers in structs should trigger an - attach action. */ - if (attach_detach - && ((region_type & (ORT_ACC | ORT_TARGET | ORT_TARGET_DATA)) - || code == OMP_TARGET_ENTER_DATA - || code == OMP_TARGET_EXIT_DATA)) - - { - /* Turn a GOMP_MAP_ATTACH_DETACH clause into a - GOMP_MAP_ATTACH or GOMP_MAP_DETACH clause after we - have detected a case that needs a GOMP_MAP_STRUCT - mapping added. */ - gomp_map_kind k - = ((code == OACC_EXIT_DATA || code == OMP_TARGET_EXIT_DATA) - ? GOMP_MAP_DETACH : GOMP_MAP_ATTACH); - OMP_CLAUSE_SET_MAP_KIND (c, k); - has_attachments = true; - } - if ((DECL_P (decl) - && (n == NULL || (n->value & GOVD_MAP) == 0)) - || (!DECL_P (decl) - && (!struct_map_to_clause - || struct_map_to_clause->get (decl) == NULL))) + bool cont = false; + tree add_decl + = build_struct_group (ctx, region_type, code, decl, pd, + component_ref_p, &flags, c, + struct_map_to_clause, prev_list_p, + list_p, pre_p, &cont); + if (add_decl == error_mark_node) { - tree l = build_omp_clause (OMP_CLAUSE_LOCATION (c), - OMP_CLAUSE_MAP); - gomp_map_kind k = attach ? GOMP_MAP_FORCE_PRESENT - : GOMP_MAP_STRUCT; - - OMP_CLAUSE_SET_MAP_KIND (l, k); - if (base_ref) - OMP_CLAUSE_DECL (l) = unshare_expr (base_ref); - else - { - OMP_CLAUSE_DECL (l) = unshare_expr (decl); - if (!DECL_P (OMP_CLAUSE_DECL (l)) - && (gimplify_expr (&OMP_CLAUSE_DECL (l), - pre_p, NULL, is_gimple_lvalue, - fb_lvalue) - == GS_ERROR)) - { - remove = true; - break; - } - } - OMP_CLAUSE_SIZE (l) - = (!attach - ? size_int (1) - : DECL_P (OMP_CLAUSE_DECL (l)) - ? DECL_SIZE_UNIT (OMP_CLAUSE_DECL (l)) - : TYPE_SIZE_UNIT (TREE_TYPE (OMP_CLAUSE_DECL (l)))); - if (struct_map_to_clause == NULL) - struct_map_to_clause - = new hash_map; - struct_map_to_clause->put (decl, l); - if (ptr || attach_detach) - { - insert_struct_comp_map (code, c, l, *prev_list_p, - NULL); - *prev_list_p = l; - prev_list_p = NULL; - } - else - { - OMP_CLAUSE_CHAIN (l) = c; - *list_p = l; - list_p = &OMP_CLAUSE_CHAIN (l); - } - if (base_ref && code == OMP_TARGET) - { - tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), - OMP_CLAUSE_MAP); - enum gomp_map_kind mkind - = GOMP_MAP_FIRSTPRIVATE_REFERENCE; - OMP_CLAUSE_SET_MAP_KIND (c2, mkind); - OMP_CLAUSE_DECL (c2) = decl; - OMP_CLAUSE_SIZE (c2) = size_zero_node; - OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (l); - OMP_CLAUSE_CHAIN (l) = c2; - } - flags = GOVD_MAP | GOVD_EXPLICIT; - if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) - || ptr - || attach_detach) - flags |= GOVD_SEEN; - if (has_attachments) - flags |= GOVD_MAP_HAS_ATTACHMENTS; - - /* If this is a *pointer-to-struct expression, make sure a - firstprivate map of the base-pointer exists. */ - if (component_ref_p - && ((TREE_CODE (decl) == MEM_REF - && integer_zerop (TREE_OPERAND (decl, 1))) - || INDIRECT_REF_P (decl)) - && DECL_P (TREE_OPERAND (decl, 0)) - && !splay_tree_lookup (ctx->variables, - ((splay_tree_key) - TREE_OPERAND (decl, 0)))) - { - decl = TREE_OPERAND (decl, 0); - tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), - OMP_CLAUSE_MAP); - enum gomp_map_kind mkind - = GOMP_MAP_FIRSTPRIVATE_POINTER; - OMP_CLAUSE_SET_MAP_KIND (c2, mkind); - OMP_CLAUSE_DECL (c2) = decl; - OMP_CLAUSE_SIZE (c2) = size_zero_node; - OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c); - OMP_CLAUSE_CHAIN (c) = c2; - } - - if (DECL_P (decl)) - goto do_add_decl; + remove = true; + break; } - else if (struct_map_to_clause) + else if (add_decl && DECL_P (add_decl)) { - tree *osc = struct_map_to_clause->get (decl); - tree *sc = NULL, *scp = NULL; - if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) - || ptr - || attach_detach) - n->value |= GOVD_SEEN; - sc = &OMP_CLAUSE_CHAIN (*osc); - if (*sc != c - && (OMP_CLAUSE_MAP_KIND (*sc) - == GOMP_MAP_FIRSTPRIVATE_REFERENCE)) - sc = &OMP_CLAUSE_CHAIN (*sc); - /* Here "prev_list_p" is the end of the inserted - alloc/release nodes after the struct node, OSC. */ - for (; *sc != c; sc = &OMP_CLAUSE_CHAIN (*sc)) - if ((ptr || attach_detach) && sc == prev_list_p) - break; - else if (TREE_CODE (OMP_CLAUSE_DECL (*sc)) - != COMPONENT_REF - && (TREE_CODE (OMP_CLAUSE_DECL (*sc)) - != INDIRECT_REF) - && (TREE_CODE (OMP_CLAUSE_DECL (*sc)) - != ARRAY_REF)) - break; - else - { - tree sc_decl = OMP_CLAUSE_DECL (*sc); - poly_offset_int offsetn; - poly_int64 bitposn; - tree base - = extract_base_bit_offset (sc_decl, NULL, - &bitposn, &offsetn); - if (base != decl) - break; - if (scp) - continue; - if ((region_type & ORT_ACC) != 0) - { - /* This duplicate checking code is currently only - enabled for OpenACC. */ - tree d1 = OMP_CLAUSE_DECL (*sc); - tree d2 = OMP_CLAUSE_DECL (c); - while (TREE_CODE (d1) == ARRAY_REF) - d1 = TREE_OPERAND (d1, 0); - while (TREE_CODE (d2) == ARRAY_REF) - d2 = TREE_OPERAND (d2, 0); - if (TREE_CODE (d1) == INDIRECT_REF) - d1 = TREE_OPERAND (d1, 0); - if (TREE_CODE (d2) == INDIRECT_REF) - d2 = TREE_OPERAND (d2, 0); - while (TREE_CODE (d1) == COMPONENT_REF) - if (TREE_CODE (d2) == COMPONENT_REF - && TREE_OPERAND (d1, 1) - == TREE_OPERAND (d2, 1)) - { - d1 = TREE_OPERAND (d1, 0); - d2 = TREE_OPERAND (d2, 0); - } - else - break; - if (d1 == d2) - { - error_at (OMP_CLAUSE_LOCATION (c), - "%qE appears more than once in map " - "clauses", OMP_CLAUSE_DECL (c)); - remove = true; - break; - } - } - if (maybe_lt (offset1, offsetn) - || (known_eq (offset1, offsetn) - && maybe_lt (bitpos1, bitposn))) - { - if (ptr || attach_detach) - scp = sc; - else - break; - } - } - if (remove) - break; - if (!attach) - OMP_CLAUSE_SIZE (*osc) - = size_binop (PLUS_EXPR, OMP_CLAUSE_SIZE (*osc), - size_one_node); - if (ptr || attach_detach) - { - tree cl = insert_struct_comp_map (code, c, NULL, - *prev_list_p, scp); - if (sc == prev_list_p) - { - *sc = cl; - prev_list_p = NULL; - } - else - { - *prev_list_p = OMP_CLAUSE_CHAIN (c); - list_p = prev_list_p; - prev_list_p = NULL; - OMP_CLAUSE_CHAIN (c) = *sc; - *sc = cl; - continue; - } - } - else if (*sc != c) - { - if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, - fb_lvalue) - == GS_ERROR) - { - remove = true; - break; - } - *list_p = OMP_CLAUSE_CHAIN (c); - OMP_CLAUSE_CHAIN (c) = *sc; - *sc = c; - continue; - } + decl = add_decl; + goto do_add_decl; } + if (cont) + continue; } else if ((code == OACC_ENTER_DATA || code == OACC_EXIT_DATA From patchwork Tue May 11 08:56:33 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1476982 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=2620:52:3:1:0:246e:9693:128c; helo=sourceware.org; envelope-from=gcc-patches-bounces@gcc.gnu.org; receiver=) Received: from 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 RSA-PSS (4096 bits) server-digest SHA256) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 4FfWyd0SSWz9sSn for ; Tue, 11 May 2021 18:57:28 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id CA9CC3886C6F; Tue, 11 May 2021 08:57:11 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa2.mentor.iphmx.com (esa2.mentor.iphmx.com [68.232.141.98]) by sourceware.org (Postfix) with ESMTPS id 64D143839C4D; Tue, 11 May 2021 08:57:08 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org 64D143839C4D Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=Julian_Brown@mentor.com IronPort-SDR: dMW1ZhkVmqyly415YPeYpzy9KvNZO/lvyNy2s31tTWdfdEaVpIDDx1BJ+6x1lMYRL0v2b64U6D MTR0Kx+66bou0925S/8orGZ/C7QUbGcS3YoD7Q1nifxMlJvw1MnS8iQuqq8PKM96tNGmFkpxE9 fpHqMgjBZZhAaGayp8i8hUIgMG0hlZFdLZG+TBg/idq3RZL6Hi66YejaLDfefRgiv8333fYUO1 ovAe0/GjJaJ6V3AYYKmvEH8v/6GFdEk4FjqV+obKtvjNYmdiRAQ3hVHQWvRYUcaqphD+SkiGUY KqI= X-IronPort-AV: E=Sophos;i="5.82,290,1613462400"; d="scan'208";a="61053251" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa2.mentor.iphmx.com with ESMTP; 11 May 2021 00:57:04 -0800 IronPort-SDR: zNmCXK1T+avSCq01Xpgb7ftWVWo+qQJPaUfGLlTOhQG4ZQeNHJPvkrVHTpPX0tYW7mT/shdXDz hKkNXq4bfIXgzwlAF/rLnX5gglee6WYyJZClLVe48Qn3CD2qCbWh7IoWDFOEQWr4yAJ5WR4hoe h0BnyWjk7e6e7gK8L01HyJEqMjkDCgK25Rtq1TQzgPPAJ4oKMptVDX1mscFTtg2F8YEQZMolEu wyUAsOJOTo3zsw6cnoynYThuaLENLQyI06vUsFGm6tCNELo6kGsxN3csua3Y+hF8e+wIktrsta rMg= From: Julian Brown To: Subject: [PATCH 3/7] [og10] Revert gimplify.c parts of "Fix template case of non-static member access inside member functions" Date: Tue, 11 May 2021 01:56:33 -0700 Message-ID: <7c0034b155be67cae409283e8338cae6ae92d4cc.1620721888.git.julian@codesourcery.com> X-Mailer: git-send-email 2.29.2 In-Reply-To: References: MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-09.mgc.mentorg.com (139.181.222.9) To SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) X-Spam-Status: No, score=-12.0 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, RCVD_IN_DNSWL_NONE, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.2 X-Spam-Checker-Version: SpamAssassin 3.4.2 (2018-09-13) 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: , Cc: Jakub Jelinek , Tobias Burnus , fortran@gcc.gnu.org, Thomas Schwinge Errors-To: gcc-patches-bounces@gcc.gnu.org Sender: "Gcc-patches" With the "rework indirect struct handling" patch later in this series, some parts of this earlier patch (by Chung-Lin) become unnecessary. This patch reverts those bits. 2021-05-11 Julian Brown gcc/ * gimplify.c (gimplify_scan_omp_clauses): Don't strip nops in indir_p case. Don't add map(*pointer_to_struct) mappings to struct_deref_set. --- gcc/gimplify.c | 19 ------------------- 1 file changed, 19 deletions(-) diff --git a/gcc/gimplify.c b/gcc/gimplify.c index ad192b27208..0674d882861 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -9574,7 +9574,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, { indir_p = true; decl = TREE_OPERAND (decl, 0); - STRIP_NOPS (decl); } if (TREE_CODE (decl) == INDIRECT_REF && DECL_P (TREE_OPERAND (decl, 0)) @@ -9747,24 +9746,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, break; } - /* If this was of the form map(*pointer_to_struct), then the - 'pointer_to_struct' DECL should be considered deref'ed. */ - if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALLOC - || GOMP_MAP_COPY_TO_P (OMP_CLAUSE_MAP_KIND (c)) - || GOMP_MAP_COPY_FROM_P (OMP_CLAUSE_MAP_KIND (c))) - && INDIRECT_REF_P (orig_decl) - && DECL_P (TREE_OPERAND (orig_decl, 0)) - && TREE_CODE (TREE_TYPE (orig_decl)) == RECORD_TYPE) - { - tree ptr = TREE_OPERAND (orig_decl, 0); - if (!struct_deref_set || !struct_deref_set->contains (ptr)) - { - if (!struct_deref_set) - struct_deref_set = new hash_set (); - struct_deref_set->add (ptr); - } - } - if (!remove && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_POINTER && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH_DETACH From patchwork Tue May 11 08:56:34 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1476983 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=2620:52:3:1:0:246e:9693:128c; helo=sourceware.org; envelope-from=gcc-patches-bounces@gcc.gnu.org; receiver=) Received: from 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 RSA-PSS (4096 bits) server-digest SHA256) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 4FfWyj5ZmWz9sT6 for ; Tue, 11 May 2021 18:57:33 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 7787F3886C58; Tue, 11 May 2021 08:57:14 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa1.mentor.iphmx.com (esa1.mentor.iphmx.com [68.232.129.153]) by sourceware.org (Postfix) with ESMTPS id 1EC353886C58; Tue, 11 May 2021 08:57:11 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org 1EC353886C58 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=Julian_Brown@mentor.com IronPort-SDR: CHBXtRdQwEUV3AkD6MAaOuCjXZaq2nq9ljW7Oz6aH4M/Qia92yqEk2fWzCTv93Cx2DsgW6+CGs I55oWlnK0vBPrrUHlyzbvYZt84kP8RGDrQVwhQsra4MRgv/AWSm8uVWFXnY12pYVA9Zt0n4cvy mdetQChSQRC2QoQj+8e1zngrULDL9uylE+DfdSkw7rYO3yzoxt+4xW/pWwXYv+oW6rIV2ottDE QWDYC2QzCUpAG/cUs9sro7QfyL6aNB3WLHsn9fUSPH59yDPp7aWhPTosiWK/pB/OYBxPtEcfxE yH8= X-IronPort-AV: E=Sophos;i="5.82,290,1613462400"; d="scan'208";a="63457170" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa1.mentor.iphmx.com with ESMTP; 11 May 2021 00:57:09 -0800 IronPort-SDR: D9alfij0yAs5VLQ1CwYkuLUNnjrksVVb6cR62lDALfkbQ1gjrGNVCaNzZ3ZFv+YEIu2Rc9WiNC ZeBsuGK2qEbYa+bH2+drUfsI+/D0g8Nb4apaaLZaJS9hJPMUS+vV6D8sRChVgw1urAnAaNGU/p dmC3oVOEHJ8B/GUndvTPZFTdHCJ65j6HniyzrRhOBIIj5wScLjLcTaYHCHae2B8Gh63d3/k2FV IJZ7SBd7lC8Ft2qhQwuyKBPLyv+sZ3WgAs3k9WG/ADquikOgwb4fRdTjrlccNx+LYFEdAReBFG OgE= From: Julian Brown To: Subject: [PATCH 4/7] [og10] Revert gimplify.c parts of "Arrow operator handling for C front-end in OpenMP map clauses" Date: Tue, 11 May 2021 01:56:34 -0700 Message-ID: <8e978e55e80508930d32e88e563279f4ce491389.1620721888.git.julian@codesourcery.com> X-Mailer: git-send-email 2.29.2 In-Reply-To: References: MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-09.mgc.mentorg.com (139.181.222.9) To SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) X-Spam-Status: No, score=-12.0 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.2 X-Spam-Checker-Version: SpamAssassin 3.4.2 (2018-09-13) 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: , Cc: Jakub Jelinek , Tobias Burnus , fortran@gcc.gnu.org, Thomas Schwinge Errors-To: gcc-patches-bounces@gcc.gnu.org Sender: "Gcc-patches" With the "rework indirect struct handling" patch later in this series, some parts of this earlier patch (by Chung-Lin) become unnecessary. This patch reverts those bits. An XFAIL has been added for a test that fails for the time being with this reversion, until the later patch in the series fixes it again. 2021-05-11 Julian Brown gcc/ * gimplify.c (build_struct_group): Remove COMPONENT_REF_P parameter. Don't call gimplify_expr on decl in non-reference case. Remove code to add FIRSTPRIVATE_POINTER for *pointer-to-struct expressions. (gimplify_scan_omp_clauses): Remove COMPONENT_REF_P handling. gcc/testsuite/ * gcc.dg/gomp/target-3.c: XFAIL test. --- gcc/gimplify.c | 41 ++++------------------------ gcc/testsuite/gcc.dg/gomp/target-3.c | 2 +- 2 files changed, 6 insertions(+), 37 deletions(-) diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 0674d882861..c2072c7188f 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -8686,8 +8686,7 @@ move_concat_nodes_after (tree first_new, tree *last_new_tail, tree *first_ptr, static tree build_struct_group (struct gimplify_omp_ctx *ctx, enum omp_region_type region_type, enum tree_code code, - tree decl, tree *pd, bool component_ref_p, - unsigned int *flags, tree c, + tree decl, tree *pd, unsigned int *flags, tree c, hash_map *&struct_map_to_clause, tree *&prev_list_p, tree *&list_p, gimple_seq *pre_p, bool *cont) @@ -8737,13 +8736,7 @@ build_struct_group (struct gimplify_omp_ctx *ctx, if (base_ref) OMP_CLAUSE_DECL (l) = unshare_expr (base_ref); else - { - OMP_CLAUSE_DECL (l) = unshare_expr (decl); - if (!DECL_P (OMP_CLAUSE_DECL (l)) - && (gimplify_expr (&OMP_CLAUSE_DECL (l), pre_p, NULL, - is_gimple_lvalue, fb_lvalue) == GS_ERROR)) - return error_mark_node; - } + OMP_CLAUSE_DECL (l) = decl; OMP_CLAUSE_SIZE (l) = (!attach ? size_int (1) : (DECL_P (OMP_CLAUSE_DECL (l)) @@ -8785,27 +8778,6 @@ build_struct_group (struct gimplify_omp_ctx *ctx, *flags |= GOVD_SEEN; if (has_attachments) *flags |= GOVD_MAP_HAS_ATTACHMENTS; - - /* If this is a *pointer-to-struct expression, make sure a - firstprivate map of the base-pointer exists. */ - if (component_ref_p - && ((TREE_CODE (decl) == MEM_REF - && integer_zerop (TREE_OPERAND (decl, 1))) - || INDIRECT_REF_P (decl)) - && DECL_P (TREE_OPERAND (decl, 0)) - && !splay_tree_lookup (ctx->variables, - ((splay_tree_key) TREE_OPERAND (decl, 0)))) - { - decl = TREE_OPERAND (decl, 0); - tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); - enum gomp_map_kind mkind = GOMP_MAP_FIRSTPRIVATE_POINTER; - OMP_CLAUSE_SET_MAP_KIND (c2, mkind); - OMP_CLAUSE_DECL (c2) = decl; - OMP_CLAUSE_SIZE (c2) = size_zero_node; - OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c); - OMP_CLAUSE_CHAIN (c) = c2; - } - return decl; } else if (struct_map_to_clause) @@ -9660,9 +9632,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH) OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALWAYS_POINTER); if ((DECL_P (decl) - || (component_ref_p - && (INDIRECT_REF_P (decl) - || TREE_CODE (decl) == MEM_REF))) + || (component_ref_p && INDIRECT_REF_P (decl))) && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_DETACH @@ -9710,9 +9680,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, bool cont = false; tree add_decl = build_struct_group (ctx, region_type, code, decl, pd, - component_ref_p, &flags, c, - struct_map_to_clause, prev_list_p, - list_p, pre_p, &cont); + &flags, c, struct_map_to_clause, + prev_list_p, list_p, pre_p, &cont); if (add_decl == error_mark_node) { remove = true; diff --git a/gcc/testsuite/gcc.dg/gomp/target-3.c b/gcc/testsuite/gcc.dg/gomp/target-3.c index 3e7921270c9..08e42eeb304 100644 --- a/gcc/testsuite/gcc.dg/gomp/target-3.c +++ b/gcc/testsuite/gcc.dg/gomp/target-3.c @@ -13,4 +13,4 @@ void foo (struct S *s) #pragma omp target enter data map (alloc: s->a, s->b) } -/* { dg-final { scan-tree-dump-times "map\\(struct:\\*s \\\[len: 2\\\]\\) map\\(alloc:s->a \\\[len: \[0-9\]+\\\]\\) map\\(alloc:s->b \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "map\\(struct:\\*s \\\[len: 2\\\]\\) map\\(alloc:s->a \\\[len: \[0-9\]+\\\]\\) map\\(alloc:s->b \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" { xfail *-*-* } } } */ From patchwork Tue May 11 08:57:44 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1476984 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=2620:52:3:1:0:246e:9693:128c; helo=sourceware.org; envelope-from=gcc-patches-bounces@gcc.gnu.org; receiver=) Received: from 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 RSA-PSS (4096 bits) server-digest SHA256) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 4FfWzH2RJyz9sjJ for ; Tue, 11 May 2021 18:58:03 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 3E6833886C51; Tue, 11 May 2021 08:58:01 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa2.mentor.iphmx.com (esa2.mentor.iphmx.com [68.232.141.98]) by sourceware.org (Postfix) with ESMTPS id 025C03865479; Tue, 11 May 2021 08:57:57 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org 025C03865479 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=Julian_Brown@mentor.com IronPort-SDR: UBuxZsLjwx4BGCvnx3r7Q86YXSkIw8/5AqpFEwjMGY1VFLIMu08HFHCiAYlo1KjuyfULdbI9SW 7Tu9AngzHP/EwzYWq49UVj1AAjQO8PrLqusg1K5eWQHVeXSBPDoYLvsauy+5xwaDjoImWyJsHQ bnmUiYQmDtkXskBwngysUubT6BXApkbb8PjNangmqU3RTIDRMGbLLYVhhRx5OWj42WqeqNokX+ O/xUDSfoZtJhl3CDrMaW/Ha+2Va0FYcpWvzBGgmAzvULTDl0uOAEaW+2PLyGtnH/2id8Wo0N3H yRk= X-IronPort-AV: E=Sophos;i="5.82,290,1613462400"; d="scan'208";a="61053287" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa2.mentor.iphmx.com with ESMTP; 11 May 2021 00:57:57 -0800 IronPort-SDR: ii62JAPFUKWZUgfZFemsFvkvtFt3Fe9/sVoixnlBPl5aBhahjGMTqwV4qFIVjq0sYDgr+lS0XO 46nYYuc9J5mbLbpXOQLGt93iQ/xGrGLkjRm1SNVy4RoAc6CetT24Z+XAgpJuAYqX0z/zccM44O DJ3PvxAyQ5RlywbAX98lkWVX2S9WF4TYs2xs4mi48QOeD62lIXwn9vzq+HdWBPR+/2dufLWnMv Qq5/PbE3lMkwqVo6NTeTx/N9RutNJDJFVF38Mkh9HpsD8lqy4VKgL0eP+eFayH3a90lTRIvH3Q 0ds= From: Julian Brown To: Subject: [PATCH 5/7] [og10] Rewrite GOMP_MAP_ATTACH_DETACH mappings for OpenMP also Date: Tue, 11 May 2021 01:57:44 -0700 Message-ID: X-Mailer: git-send-email 2.29.2 In-Reply-To: References: MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: SVR-IES-MBX-03.mgc.mentorg.com (139.181.222.3) To SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) X-Spam-Status: No, score=-12.0 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, RCVD_IN_DNSWL_NONE, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.2 X-Spam-Checker-Version: SpamAssassin 3.4.2 (2018-09-13) 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: , Cc: Jakub Jelinek , Tobias Burnus , fortran@gcc.gnu.org, Thomas Schwinge Errors-To: gcc-patches-bounces@gcc.gnu.org Sender: "Gcc-patches" It never makes sense for a GOMP_MAP_ATTACH_DETACH mapping to survive beyond gimplify.c, and with OpenMP making use of that mapping type too now alongside OpenACC, there are cases where it was making it through to omp-low.c. This patch rewrites such mappings to GOMP_MAP_ATTACH or GOMP_MAP_DETACH unconditionally for both OpenACC and OpenMP, in cases where it hasn't otherwise been handled already in the preceding code. 2021-05-11 Julian Brown gcc/ * gimplify.c (gimplify_scan_omp_clauses): Remove OpenACC-only condition for changing GOMP_MAP_ATTACH_DETACH to GOMP_MAP_ATTACH or GOMP_MAP_DETACH. Use detach for OMP_TARGET_EXIT_DATA also. --- gcc/gimplify.c | 13 ++++--------- 1 file changed, 4 insertions(+), 9 deletions(-) diff --git a/gcc/gimplify.c b/gcc/gimplify.c index c2072c7188f..86000f8470b 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -9695,16 +9695,11 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, if (cont) continue; } - else if ((code == OACC_ENTER_DATA - || code == OACC_EXIT_DATA - || code == OACC_DATA - || code == OACC_PARALLEL - || code == OACC_KERNELS - || code == OACC_SERIAL) - && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH) + else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH) { - gomp_map_kind k = (code == OACC_EXIT_DATA - ? GOMP_MAP_DETACH : GOMP_MAP_ATTACH); + gomp_map_kind k + = ((code == OACC_EXIT_DATA || code == OMP_TARGET_EXIT_DATA) + ? GOMP_MAP_DETACH : GOMP_MAP_ATTACH); OMP_CLAUSE_SET_MAP_KIND (c, k); } From patchwork Tue May 11 08:57:45 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1476985 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=8.43.85.97; helo=sourceware.org; envelope-from=gcc-patches-bounces@gcc.gnu.org; receiver=) Received: from sourceware.org (ip-8-43-85-97.sourceware.org [8.43.85.97]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature RSA-PSS (4096 bits) server-digest SHA256) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 4FfWzR4DH8z9sSn for ; Tue, 11 May 2021 18:58:11 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 340263886C75; Tue, 11 May 2021 08:58:04 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa2.mentor.iphmx.com (esa2.mentor.iphmx.com [68.232.141.98]) by sourceware.org (Postfix) with ESMTPS id 0BD983839C4D; Tue, 11 May 2021 08:57:58 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org 0BD983839C4D Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=Julian_Brown@mentor.com IronPort-SDR: TteNtjgcRaOlh6tKcMRIa39ajgYj5K7HujGQHb5RX5WCA9/gDzXAQwLTTX0uCF9OetdT9DCLQL g7xxHWdJZr5Dp/rOQ1W0y1N+1xaBFvuapY4+U7nRsunH5VQUEethCxP+m1TZrgUBwY1d7F2Xhb i6Xudj7II8xBIAYomsHmIZ3Hx2qTqmJjhczOUcFadAiwdRwQSb7CcZBecdJIOzb4kZ//s2MaWk CZrshv1FhVzpZh0fIVsTwEBbn/IfpeBRCEvVWus1jJyAIEf8YyInb/zu3tZmrhCPxgVpmjPq8m USo= X-IronPort-AV: E=Sophos;i="5.82,290,1613462400"; d="scan'208";a="61053289" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa2.mentor.iphmx.com with ESMTP; 11 May 2021 00:57:58 -0800 IronPort-SDR: 6a37Koz/rQUhbj/E3WgbuSrKchIUbCJCN5N/I6ZdqPO4O2o4P8Nad5RwRa0g2JUoIdgipNzG2B L8waKSUZwnMSRNq/mtQRYGGxCiLssgyTG1fGfj0x2uTDBpv9k3jjhCb2lElq67HFseRcYriYut j9+AOQjJjUdaayv48rCAxFUq9KlVTEyAIq8rr5adEPP+39Iv7ta+tPf4+enb3HzeqMxq0TZf+z K58CSrpB6nnYnJPKvZt7/S/tHshzCfu5lkoWfKIru+gQ64UiW6jjPBW0cY++4AUNLMvcZKbhvk CLg= From: Julian Brown To: Subject: [PATCH 6/7] [og10] Rework indirect struct handling for OpenACC/OpenMP in gimplify.c Date: Tue, 11 May 2021 01:57:45 -0700 Message-ID: <38a7e91eab4d93dd4c8021d1aa19ac5cb5d67264.1620721888.git.julian@codesourcery.com> X-Mailer: git-send-email 2.29.2 In-Reply-To: References: MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: SVR-IES-MBX-03.mgc.mentorg.com (139.181.222.3) To SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) X-Spam-Status: No, score=-12.0 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, RCVD_IN_DNSWL_NONE, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.2 X-Spam-Checker-Version: SpamAssassin 3.4.2 (2018-09-13) 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: , Cc: Jakub Jelinek , Tobias Burnus , fortran@gcc.gnu.org, Thomas Schwinge Errors-To: gcc-patches-bounces@gcc.gnu.org Sender: "Gcc-patches" This patch reworks indirect struct handling in gimplify.c (i.e. for struct components mapped with "mystruct->a[0:n]", "mystruct->b", etc.), for both OpenACC and OpenMP. The key observation leading to these changes was that component mappings of references-to-structures is already implemented and working, and indirect struct component handling via a pointer can work quite similarly. That lets us remove some earlier, special-case handling for mapping indirect struct component accesses for OpenACC, which required the pointed-to struct to be manually mapped before the indirect component mapping. With this patch, you can map struct components directly (e.g. an array slice "mystruct->a[0:n]") just like you can map a non-indirect struct component slice ("mystruct.a[0:n]"). Both references-to-pointers (with the former syntax) and references to structs (with the latter syntax) work now. For Fortran class pointers, we no longer re-use GOMP_MAP_TO_PSET for the class metadata (the structure that points to the class data and vptr) -- it is instead treated as any other struct. For C++, the struct handling also works for class members ("this->foo"), without having to explicitly map "this[:1]" first. For OpenACC, we permit chained indirect component references ("mystruct->a->b[0:n]"), though only the last part of such mappings will trigger an attach/detach operation. To properly use such a construct on the target, you must still manually map "mystruct->a[:1]" first -- but there's no need to map "mystruct[:1]" explicitly before that. 2021-05-11 Julian Brown gcc/fortran/ * trans-openmp.c (gfc_trans_omp_clauses): Don't create GOMP_MAP_TO_PSET mappings for class metadata, nor GOMP_MAP_POINTER mappings for POINTER_TYPE_P decls. gcc/ * gimplify.c (extract_base_bit_offset): Add BASE_IND parameter. Handle pointer-typed indirect references alongside reference-typed ones. (strip_components, strip_components_and_deref, aggregate_base_p): New functions. (build_struct_group): Remove PD parameter. Add pointer type indirect ref handling, including chained references. Handle pointers and references to structs in OpenACC regions as well as OpenMP ones. Remove gimplification of non-pointer struct component mappings. (gimplify_scan_omp_clauses): Remove struct_deref_set handling. Rework pointer-type indirect structure access handling to work more like the reference-typed handling. * omp-low.c (scan_sharing_clauses): Handle pointer-type indirect struct references, and references to pointers to structs also. gcc/testsuite/ * g++.dg/goacc/member-array-acc.C: New test. * g++.dg/gomp/member-array-omp.C: New test. * g++.dg/gomp/target-3.C: Adjust scan dump matching patterns. * g++.dg/gomp/target-this-2.C: Adjust scan dump matching patterns. * gcc.dg/gomp/target-3.c: Remove XFAIL. libgomp/ * testsuite/libgomp.oacc-c-c++-common/deep-copy-15.c: New test. * testsuite/libgomp.oacc-c-c++-common/deep-copy-16.c: New test. --- gcc/fortran/trans-openmp.c | 20 +- gcc/gimplify.c | 321 +++++++++++------- gcc/omp-low.c | 16 +- gcc/testsuite/g++.dg/goacc/member-array-acc.C | 13 + gcc/testsuite/g++.dg/gomp/member-array-omp.C | 13 + gcc/testsuite/g++.dg/gomp/target-3.C | 4 +- gcc/testsuite/g++.dg/gomp/target-this-2.C | 2 +- gcc/testsuite/gcc.dg/gomp/target-3.c | 2 +- .../libgomp.oacc-c-c++-common/deep-copy-15.c | 68 ++++ .../libgomp.oacc-c-c++-common/deep-copy-16.c | 95 ++++++ 10 files changed, 405 insertions(+), 149 deletions(-) create mode 100644 gcc/testsuite/g++.dg/goacc/member-array-acc.C create mode 100644 gcc/testsuite/g++.dg/gomp/member-array-omp.C create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-15.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-16.c diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index e3df4bbf84e..9098b35c9f1 100644 --- a/gcc/fortran/trans-openmp.c +++ b/gcc/fortran/trans-openmp.c @@ -2693,30 +2693,16 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, tree present = gfc_omp_check_optional_argument (decl, true); if (openacc && n->sym->ts.type == BT_CLASS) { - tree type = TREE_TYPE (decl); if (n->sym->attr.optional) sorry ("optional class parameter"); - if (POINTER_TYPE_P (type)) - { - node4 = build_omp_clause (input_location, - OMP_CLAUSE_MAP); - OMP_CLAUSE_SET_MAP_KIND (node4, GOMP_MAP_POINTER); - OMP_CLAUSE_DECL (node4) = decl; - OMP_CLAUSE_SIZE (node4) = size_int (0); - decl = build_fold_indirect_ref (decl); - } tree ptr = gfc_class_data_get (decl); ptr = build_fold_indirect_ref (ptr); OMP_CLAUSE_DECL (node) = ptr; OMP_CLAUSE_SIZE (node) = gfc_class_vtab_size_get (decl); node2 = build_omp_clause (input_location, OMP_CLAUSE_MAP); - OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_TO_PSET); - OMP_CLAUSE_DECL (node2) = decl; - OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type); - node3 = build_omp_clause (input_location, OMP_CLAUSE_MAP); - OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_ATTACH_DETACH); - OMP_CLAUSE_DECL (node3) = gfc_class_data_get (decl); - OMP_CLAUSE_SIZE (node3) = size_int (0); + OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_ATTACH_DETACH); + OMP_CLAUSE_DECL (node2) = gfc_class_data_get (decl); + OMP_CLAUSE_SIZE (node2) = size_int (0); goto finalize_map_clause; } else if (POINTER_TYPE_P (TREE_TYPE (decl)) diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 86000f8470b..6d204908c82 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -8323,8 +8323,8 @@ build_struct_comp_nodes (enum tree_code code, tree grp_start, tree grp_end, has array type, else return NULL. */ static tree -extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp, - poly_offset_int *poffsetp) +extract_base_bit_offset (tree base, tree *base_ind, tree *base_ref, + poly_int64 *bitposp, poly_offset_int *poffsetp) { tree offset; poly_int64 bitsize, bitpos; @@ -8332,6 +8332,9 @@ extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp, int unsignedp, reversep, volatilep = 0; poly_offset_int poffset; + if (base_ind) + *base_ind = NULL_TREE; + if (base_ref) *base_ref = NULL_TREE; @@ -8352,14 +8355,27 @@ extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp, base = get_inner_reference (base, &bitsize, &bitpos, &offset, &mode, &unsignedp, &reversep, &volatilep); - tree orig_base = base; - + if ((TREE_CODE (base) == INDIRECT_REF + || (TREE_CODE (base) == MEM_REF + && integer_zerop (TREE_OPERAND (base, 1)))) + && TREE_CODE (TREE_TYPE (TREE_OPERAND (base, 0))) == POINTER_TYPE) + { + if (base_ind) + *base_ind = base; + base = TREE_OPERAND (base, 0); + } if ((TREE_CODE (base) == INDIRECT_REF || (TREE_CODE (base) == MEM_REF && integer_zerop (TREE_OPERAND (base, 1)))) && DECL_P (TREE_OPERAND (base, 0)) && TREE_CODE (TREE_TYPE (TREE_OPERAND (base, 0))) == REFERENCE_TYPE) - base = TREE_OPERAND (base, 0); + { + if (base_ref) + *base_ref = base; + base = TREE_OPERAND (base, 0); + } + + STRIP_NOPS (base); gcc_assert (offset == NULL_TREE || poly_int_tree_p (offset)); @@ -8374,10 +8390,6 @@ extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp, *bitposp = bitpos; *poffsetp = poffset; - /* Set *BASE_REF if BASE was a dereferenced reference variable. */ - if (base_ref && orig_base != base) - *base_ref = orig_base; - return base; } @@ -8394,6 +8406,59 @@ is_or_contains_p (tree expr, tree base_ptr) return expr == base_ptr; } +/* Remove COMPONENT_REFS from EXPR. */ + +static tree +strip_components (tree expr) +{ + while (TREE_CODE (expr) == COMPONENT_REF) + expr = TREE_OPERAND (expr, 0); + + return expr; +} + +/* Remove COMPONENT_REFS and indirections from EXPR. */ + +static tree +strip_components_and_deref (tree expr) +{ + while (TREE_CODE (expr) == COMPONENT_REF + || TREE_CODE (expr) == INDIRECT_REF + || (TREE_CODE (expr) == MEM_REF + && integer_zerop (TREE_OPERAND (expr, 1)))) + expr = TREE_OPERAND (expr, 0); + + return expr; +} + +/* Return TRUE if EXPR is something we will use as the base of an aggregate + access, either: + + - a DECL_P. + - a struct component with no indirection ("a.b.c"). + - a struct component with indirection ("a->b->c"). +*/ + +static bool +aggregate_base_p (tree expr) +{ + while (TREE_CODE (expr) == COMPONENT_REF + && (DECL_P (TREE_OPERAND (expr, 0)) + || (TREE_CODE (TREE_OPERAND (expr, 0)) == COMPONENT_REF))) + expr = TREE_OPERAND (expr, 0); + + if (DECL_P (expr)) + return true; + + if (TREE_CODE (expr) == COMPONENT_REF + && (TREE_CODE (TREE_OPERAND (expr, 0)) == INDIRECT_REF + || (TREE_CODE (TREE_OPERAND (expr, 0)) == MEM_REF + && integer_zerop (TREE_OPERAND (TREE_OPERAND (expr, 0), 1))))) + return true; + + return false; +} + /* Implement OpenMP 5.x map ordering rules for target directives. There are several rules, and with some level of ambiguity, hopefully we can at least collect the complexity here in one place. */ @@ -8686,20 +8751,24 @@ move_concat_nodes_after (tree first_new, tree *last_new_tail, tree *first_ptr, static tree build_struct_group (struct gimplify_omp_ctx *ctx, enum omp_region_type region_type, enum tree_code code, - tree decl, tree *pd, unsigned int *flags, tree c, + tree decl, unsigned int *flags, tree c, hash_map *&struct_map_to_clause, tree *&prev_list_p, tree *&list_p, gimple_seq *pre_p, bool *cont) { poly_offset_int coffset; poly_int64 cbitpos; - tree base_ref; + tree base_ind, base_ref; + tree *list_in_p = list_p, *prev_list_in_p = prev_list_p; - tree base = extract_base_bit_offset (OMP_CLAUSE_DECL (c), &base_ref, - &cbitpos, &coffset); + tree base = extract_base_bit_offset (OMP_CLAUSE_DECL (c), &base_ind, + &base_ref, &cbitpos, &coffset); gcc_assert (base == decl); + /* Here, DECL is usually a DECL_P, unless we have chained indirect member + accesses, e.g. mystruct->a->b. In that case it'll be the "mystruct->a" + part. */ splay_tree_node n = (DECL_P (decl) ? splay_tree_lookup (ctx->variables, (splay_tree_key) decl) @@ -8733,7 +8802,9 @@ build_struct_group (struct gimplify_omp_ctx *ctx, gomp_map_kind k = attach ? GOMP_MAP_FORCE_PRESENT : GOMP_MAP_STRUCT; OMP_CLAUSE_SET_MAP_KIND (l, k); - if (base_ref) + if (base_ind) + OMP_CLAUSE_DECL (l) = unshare_expr (base_ind); + else if (base_ref) OMP_CLAUSE_DECL (l) = unshare_expr (base_ref); else OMP_CLAUSE_DECL (l) = decl; @@ -8763,15 +8834,81 @@ build_struct_group (struct gimplify_omp_ctx *ctx, } else list_p = insert_node_after (l, list_p); - if (base_ref && code == OMP_TARGET) + + /* Handle pointers to structs and references to structs: these cases + have an additional GOMP_MAP_FIRSTPRIVATE_{REFERENCE,POINTER} node + inserted after the GOMP_MAP_STRUCT node. References to pointers + use GOMP_MAP_FIRSTPRIVATE_REFERENCE. */ + if ((base_ind || base_ref) && (region_type & ORT_TARGET)) { tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); - enum gomp_map_kind mkind = GOMP_MAP_FIRSTPRIVATE_REFERENCE; + enum gomp_map_kind mkind = base_ref ? GOMP_MAP_FIRSTPRIVATE_REFERENCE + : GOMP_MAP_FIRSTPRIVATE_POINTER; OMP_CLAUSE_SET_MAP_KIND (c2, mkind); - OMP_CLAUSE_DECL (c2) = decl; OMP_CLAUSE_SIZE (c2) = size_zero_node; - OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (l); - OMP_CLAUSE_CHAIN (l) = c2; + tree sdecl = strip_components (decl); + if (DECL_P (decl) + && (POINTER_TYPE_P (TREE_TYPE (sdecl)) + || TREE_CODE (TREE_TYPE (sdecl)) == REFERENCE_TYPE)) + { + /* Insert after struct node. */ + OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (l); + OMP_CLAUSE_DECL (c2) = decl; + OMP_CLAUSE_CHAIN (l) = c2; + } + else if (!POINTER_TYPE_P (TREE_TYPE (sdecl)) + && TREE_CODE (TREE_TYPE (sdecl)) != REFERENCE_TYPE) + { + /* If the ultimate base for this component access is not a + pointer or reference, that means it is a struct component + access itself. Insert a node to be processed on the next + iteration of our caller's loop, which will subsequently be + turned into a new GOMP_MAP_STRUCT mapping itself. + + We need to do this else the non-DECL_P base won't be + rewritten correctly in the offloaded region. */ + tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), + OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FORCE_PRESENT); + OMP_CLAUSE_DECL (c2) = unshare_expr (decl); + OMP_CLAUSE_SIZE (c2) = (DECL_P (decl) + ? DECL_SIZE_UNIT (decl) + : TYPE_SIZE_UNIT (TREE_TYPE (decl))); + tree *next_node = &OMP_CLAUSE_CHAIN (*list_p); + OMP_CLAUSE_CHAIN (c2) = *next_node; + *next_node = c2; + *cont = true; + return NULL_TREE; + } + else + { + /* We have chained indirect structure member accesses, e.g. + "mystruct->a->b". DECL isn't a DECL_P in that case (it will + be "mystruct->a"), and we can't use + GOMP_MAP_FIRSTPRIVATE_{REFERENCE,POINTER} directly after the + GOMP_MAP_STRUCT. Create a new FORCE_PRESENT mapping instead + for the ultimate root node (i.e. "mystruct"), and use + that. */ + tree use_decl = unshare_expr (sdecl); + decl = strip_components_and_deref (decl); + OMP_CLAUSE_DECL (c2) = decl; + tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c), + OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c3, GOMP_MAP_FORCE_PRESENT); + OMP_CLAUSE_SIZE (c3) = size_one_node; + OMP_CLAUSE_DECL (c3) = use_decl; + /* Insert FORCE_PRESENT and FIRSTPRIVATE_{REFERENCE,POINTER} + nodes before struct node. */ + OMP_CLAUSE_CHAIN (c3) = c2; + /* If we have a pointer mapping group, insert these nodes before + the start of that group, else insert before incoming list_p + (the currently-scanned node). */ + tree *insert_at + = (ptr || attach_detach) ? prev_list_in_p : list_in_p; + OMP_CLAUSE_CHAIN (c2) = *insert_at; + *insert_at = c3; + list_p = &OMP_CLAUSE_CHAIN (c2); + } } *flags = GOVD_MAP | GOVD_EXPLICIT; if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) || ptr || attach_detach) @@ -8784,11 +8921,18 @@ build_struct_group (struct gimplify_omp_ctx *ctx, { tree *osc = struct_map_to_clause->get (decl); tree *sc = NULL, *scp = NULL; - if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) || ptr || attach_detach) + if (n && (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) + || ptr + || attach_detach)) n->value |= GOVD_SEEN; sc = &OMP_CLAUSE_CHAIN (*osc); + /* The struct mapping might be immediately followed by a + FIRSTPRIVATE_POINTER and/or FIRSTPRIVATE_REFERENCE -- if it's an + indirect access or a reference, or both. (This added node is removed + in omp-low.c after it has been processed.) */ if (*sc != c - && OMP_CLAUSE_MAP_KIND (*sc) == GOMP_MAP_FIRSTPRIVATE_REFERENCE) + && (OMP_CLAUSE_MAP_KIND (*sc) == GOMP_MAP_FIRSTPRIVATE_POINTER + || OMP_CLAUSE_MAP_KIND (*sc) == GOMP_MAP_FIRSTPRIVATE_REFERENCE)) sc = &OMP_CLAUSE_CHAIN (*sc); for (; *sc != c; sc = &OMP_CLAUSE_CHAIN (*sc)) { @@ -8803,8 +8947,8 @@ build_struct_group (struct gimplify_omp_ctx *ctx, tree sc_decl = OMP_CLAUSE_DECL (*sc); poly_offset_int offset; poly_int64 bitpos; - tree base = extract_base_bit_offset (sc_decl, NULL, &bitpos, - &offset); + tree base = extract_base_bit_offset (sc_decl, NULL, NULL, + &bitpos, &offset); if (base != decl) break; if (scp) @@ -8927,9 +9071,6 @@ build_struct_group (struct gimplify_omp_ctx *ctx, } else if (*sc != c) { - if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, fb_lvalue) - == GS_ERROR) - return error_mark_node; /* In the non-pointer case, the mapping clause itself is moved into the correct position in the struct component list, which in this case is just SC. */ @@ -8951,7 +9092,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, struct gimplify_omp_ctx *ctx, *outer_ctx; tree c; hash_map *struct_map_to_clause = NULL; - hash_set *struct_deref_set = NULL; tree *prev_list_p = NULL, *orig_list_p = list_p; int handled_depend_iterators = -1; int nowait = -1; @@ -9518,111 +9658,43 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, && TREE_CODE (decl) == INDIRECT_REF && TREE_CODE (TREE_OPERAND (decl, 0)) == COMPONENT_REF && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0))) - == REFERENCE_TYPE) - && (OMP_CLAUSE_MAP_KIND (c) - != GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION)) + == REFERENCE_TYPE)) { pd = &TREE_OPERAND (decl, 0); decl = TREE_OPERAND (decl, 0); } - bool indir_p = false; - bool component_ref_p = false; - tree orig_decl = decl; - tree decl_ref = NULL_TREE; - if ((region_type & (ORT_ACC | ORT_TARGET | ORT_TARGET_DATA)) != 0 - && TREE_CODE (*pd) == COMPONENT_REF - && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH - && code != OACC_UPDATE) + if (TREE_CODE (decl) == COMPONENT_REF) { - while (TREE_CODE (decl) == COMPONENT_REF) - { - decl = TREE_OPERAND (decl, 0); - component_ref_p = true; - if (((TREE_CODE (decl) == MEM_REF - && integer_zerop (TREE_OPERAND (decl, 1))) - || INDIRECT_REF_P (decl)) - && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0))) - == POINTER_TYPE)) - { - indir_p = true; - decl = TREE_OPERAND (decl, 0); - } - if (TREE_CODE (decl) == INDIRECT_REF - && DECL_P (TREE_OPERAND (decl, 0)) - && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0))) - == REFERENCE_TYPE)) - { - decl_ref = decl; - decl = TREE_OPERAND (decl, 0); - } - } - } - else if (TREE_CODE (decl) == COMPONENT_REF - && (OMP_CLAUSE_MAP_KIND (c) - != GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION)) - { - component_ref_p = true; - while (TREE_CODE (decl) == COMPONENT_REF) + /* Strip off component refs from RHS of e.g. "a->b->c.d.e" + (which would leave "a->b" in that case). This is intended + to be equivalent to the base finding done by + get_inner_reference. */ + while (TREE_CODE (decl) == COMPONENT_REF + && (DECL_P (TREE_OPERAND (decl, 0)) + || (TREE_CODE (TREE_OPERAND (decl, 0)) + == COMPONENT_REF))) + decl = TREE_OPERAND (decl, 0); + + if (TREE_CODE (decl) == COMPONENT_REF) decl = TREE_OPERAND (decl, 0); + + /* Strip off RHS from "a->b". */ + if ((TREE_CODE (decl) == INDIRECT_REF + || (TREE_CODE (decl) == MEM_REF + && integer_zerop (TREE_OPERAND (decl, 1)))) + && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0))) + == POINTER_TYPE)) + decl = TREE_OPERAND (decl, 0); + + /* Strip off RHS from "a_ref.b" (where a_ref is + reference-typed). */ if (TREE_CODE (decl) == INDIRECT_REF && DECL_P (TREE_OPERAND (decl, 0)) && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0))) == REFERENCE_TYPE)) decl = TREE_OPERAND (decl, 0); - } - if (decl != orig_decl && DECL_P (decl) && indir_p) - { - gomp_map_kind k - = ((code == OACC_EXIT_DATA || code == OMP_TARGET_EXIT_DATA) - ? GOMP_MAP_DETACH : GOMP_MAP_ATTACH); - /* We have a dereference of a struct member. Make this an - attach/detach operation, and ensure the base pointer is - mapped as a FIRSTPRIVATE_POINTER. */ - OMP_CLAUSE_SET_MAP_KIND (c, k); - flags = GOVD_MAP | GOVD_SEEN | GOVD_EXPLICIT; - tree next_clause = OMP_CLAUSE_CHAIN (c); - if (k == GOMP_MAP_ATTACH - && code != OACC_ENTER_DATA - && code != OMP_TARGET_ENTER_DATA - && (!next_clause - || (OMP_CLAUSE_CODE (next_clause) != OMP_CLAUSE_MAP) - || (OMP_CLAUSE_MAP_KIND (next_clause) - != GOMP_MAP_POINTER) - || OMP_CLAUSE_DECL (next_clause) != decl) - && (!struct_deref_set - || !struct_deref_set->contains (decl))) - { - if (!struct_deref_set) - struct_deref_set = new hash_set (); - /* As well as the attach, we also need a - FIRSTPRIVATE_POINTER clause to properly map the - pointer to the struct base. */ - tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), - OMP_CLAUSE_MAP); - OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALLOC); - OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c2) - = 1; - tree charptr_zero - = build_int_cst (build_pointer_type (char_type_node), - 0); - OMP_CLAUSE_DECL (c2) - = build2 (MEM_REF, char_type_node, - decl_ref ? decl_ref : decl, charptr_zero); - OMP_CLAUSE_SIZE (c2) = size_zero_node; - tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c), - OMP_CLAUSE_MAP); - OMP_CLAUSE_SET_MAP_KIND (c3, - GOMP_MAP_FIRSTPRIVATE_POINTER); - OMP_CLAUSE_DECL (c3) = decl; - OMP_CLAUSE_SIZE (c3) = size_zero_node; - tree mapgrp = *prev_list_p; - *prev_list_p = c2; - OMP_CLAUSE_CHAIN (c3) = mapgrp; - OMP_CLAUSE_CHAIN (c2) = c3; - - struct_deref_set->add (decl); - } - goto do_add_decl; + + STRIP_NOPS (decl); } /* An "attach/detach" operation on an update directive should behave as a GOMP_MAP_ALWAYS_POINTER. Beware that @@ -9631,8 +9703,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, if (code == OACC_UPDATE && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH) OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALWAYS_POINTER); - if ((DECL_P (decl) - || (component_ref_p && INDIRECT_REF_P (decl))) + if (aggregate_base_p (decl) && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_DETACH @@ -9679,7 +9750,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, } bool cont = false; tree add_decl - = build_struct_group (ctx, region_type, code, decl, pd, + = build_struct_group (ctx, region_type, code, decl, &flags, c, struct_map_to_clause, prev_list_p, list_p, pre_p, &cont); if (add_decl == error_mark_node) @@ -10310,8 +10381,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, gimplify_omp_ctxp = ctx; if (struct_map_to_clause) delete struct_map_to_clause; - if (struct_deref_set) - delete struct_deref_set; } /* Return true if DECL is a candidate for shared to firstprivate diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 64b7c19ff39..bb473d58f12 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -1651,8 +1651,10 @@ scan_sharing_clauses (tree clauses, omp_context *ctx, if (TREE_CODE (decl) == COMPONENT_REF || (TREE_CODE (decl) == INDIRECT_REF && TREE_CODE (TREE_OPERAND (decl, 0)) == COMPONENT_REF - && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0))) - == REFERENCE_TYPE))) + && (((TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0))) + == REFERENCE_TYPE) + || (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0))) + == POINTER_TYPE))))) break; if (DECL_SIZE (decl) && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST) @@ -13504,6 +13506,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_REFERENCE) is_ref = false; bool ref_to_array = false; + bool ref_to_ptr = false; if (is_ref) { type = TREE_TYPE (type); @@ -13522,6 +13525,12 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) new_var = decl2; type = TREE_TYPE (new_var); } + else if (TREE_CODE (type) == REFERENCE_TYPE + && TREE_CODE (TREE_TYPE (type)) == POINTER_TYPE) + { + type = TREE_TYPE (type); + ref_to_ptr = true; + } x = build_receiver_ref (prev, false, ctx); x = fold_convert_loc (clause_loc, type, x); if (!integer_zerop (OMP_CLAUSE_SIZE (c))) @@ -13544,7 +13553,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 (is_ref && !ref_to_array) + if ((is_ref && !ref_to_array) + || ref_to_ptr) { tree t = create_tmp_var_raw (type, get_name (var)); gimple_add_tmp_var (t); diff --git a/gcc/testsuite/g++.dg/goacc/member-array-acc.C b/gcc/testsuite/g++.dg/goacc/member-array-acc.C new file mode 100644 index 00000000000..e0c11570f5d --- /dev/null +++ b/gcc/testsuite/g++.dg/goacc/member-array-acc.C @@ -0,0 +1,13 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +struct Foo { + float *a; + void init(int N) { + a = new float[N]; + #pragma acc enter data create(a[0:N]) + } +}; +int main() { Foo x; x.init(1024); } + +/* { dg-final { scan-tree-dump {struct:\*\(struct Foo \*\) this \[len: 1\]\) map\(alloc:\(\(struct Foo \*\) this\)->a \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: _[0-9]+\]\) map\(attach:\(\(struct Foo \*\) this\)->a \[bias: 0\]\)} "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/member-array-omp.C b/gcc/testsuite/g++.dg/gomp/member-array-omp.C new file mode 100644 index 00000000000..40fc503f882 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/member-array-omp.C @@ -0,0 +1,13 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +struct Foo { + float *a; + void init(int N) { + a = new float[N]; + #pragma omp target enter data map(alloc:a[0:N]) + } +}; +int main() { Foo x; x.init(1024); } + +/* { dg-final { scan-tree-dump {struct:\*\(struct Foo \*\) this \[len: 1\]\) map\(alloc:\(\(struct Foo \*\) this\)->a \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: _[0-9]+\]\) map\(attach:\(\(struct Foo \*\) this\)->a \[bias: 0\]\)} "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/target-3.C b/gcc/testsuite/g++.dg/gomp/target-3.C index f4d40ec8e4b..9fc0e4173d2 100644 --- a/gcc/testsuite/g++.dg/gomp/target-3.C +++ b/gcc/testsuite/g++.dg/gomp/target-3.C @@ -15,6 +15,8 @@ S::bar (int x) #pragma omp target enter data map (alloc: a, b) } +/* { dg-final { scan-tree-dump-times "map\\(struct:\\*\\(struct S \\*\\) this \\\[len: 2\\\]\\) map\\(alloc:\\(\\(struct S \\*\\) this\\)->a \\\[len: \[0-9\]+\\\]\\) map\\(alloc:\\(\\(struct S \\*\\) this\\)->b \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */ + template struct T { @@ -33,4 +35,4 @@ T::bar (int x) template struct T<0>; -/* { dg-final { scan-tree-dump-times "map\\(struct:\\*this \\\[len: 2\\\]\\) map\\(alloc:this->a \\\[len: \[0-9\]+\\\]\\) map\\(alloc:this->b \\\[len: \[0-9\]+\\\]\\)" 4 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "map\\(struct:\\*\\(struct T \\*\\) this \\\[len: 2\\\]\\) map\\(alloc:\\(\\(struct T \\*\\) this\\)->a \\\[len: \[0-9\]+\\\]\\) map\\(alloc:\\(\\(struct T \\*\\) this\\)->b \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/target-this-2.C b/gcc/testsuite/g++.dg/gomp/target-this-2.C index 679c85a54dd..a5e832130fb 100644 --- a/gcc/testsuite/g++.dg/gomp/target-this-2.C +++ b/gcc/testsuite/g++.dg/gomp/target-this-2.C @@ -46,4 +46,4 @@ int main (void) return 0; } -/* { dg-final { scan-tree-dump {map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\)} "gimple" } } */ +/* { dg-final { scan-tree-dump {map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\)} "gimple" } } */ diff --git a/gcc/testsuite/gcc.dg/gomp/target-3.c b/gcc/testsuite/gcc.dg/gomp/target-3.c index 08e42eeb304..3e7921270c9 100644 --- a/gcc/testsuite/gcc.dg/gomp/target-3.c +++ b/gcc/testsuite/gcc.dg/gomp/target-3.c @@ -13,4 +13,4 @@ void foo (struct S *s) #pragma omp target enter data map (alloc: s->a, s->b) } -/* { dg-final { scan-tree-dump-times "map\\(struct:\\*s \\\[len: 2\\\]\\) map\\(alloc:s->a \\\[len: \[0-9\]+\\\]\\) map\\(alloc:s->b \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" { xfail *-*-* } } } */ +/* { dg-final { scan-tree-dump-times "map\\(struct:\\*s \\\[len: 2\\\]\\) map\\(alloc:s->a \\\[len: \[0-9\]+\\\]\\) map\\(alloc:s->b \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-15.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-15.c new file mode 100644 index 00000000000..27fe1a9d07d --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-15.c @@ -0,0 +1,68 @@ +#include + +/* Test multiple struct dereferences on one directive, and slices starting at + non-zero. */ + +typedef struct { + int *a; + int *b; + int *c; +} mystruct; + +int main(int argc, char* argv[]) +{ + const int N = 1024; + mystruct *m = (mystruct *) malloc (sizeof (*m)); + int i; + + m->a = (int *) malloc (N * sizeof (int)); + m->b = (int *) malloc (N * sizeof (int)); + m->c = (int *) malloc (N * sizeof (int)); + + for (i = 0; i < N; i++) + { + m->a[i] = 0; + m->b[i] = 0; + m->c[i] = 0; + } + + for (int i = 0; i < 99; i++) + { + int j; +#pragma acc parallel loop copy(m->a[0:N]) + for (j = 0; j < N; j++) + m->a[j]++; +#pragma acc parallel loop copy(m->b[0:N], m->c[5:N-10]) + for (j = 0; j < N; j++) + { + m->b[j]++; + if (j > 5 && j < N - 5) + m->c[j]++; + } + } + + for (i = 0; i < N; i++) + { + if (m->a[i] != 99) + abort (); + if (m->b[i] != 99) + abort (); + if (i > 5 && i < N-5) + { + if (m->c[i] != 99) + abort (); + } + else + { + if (m->c[i] != 0) + abort (); + } + } + + free (m->a); + free (m->b); + free (m->c); + free (m); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-16.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-16.c new file mode 100644 index 00000000000..f62e190655d --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-16.c @@ -0,0 +1,95 @@ +#include + +/* Test sorting of nested struct accesses (unfinished!). */ + +typedef struct { + int *a; + int b; + int *c; +} str1; + +typedef struct { + int d; + int *e; + str1 *f; +} str2; + +typedef struct { + int g; + int h; + str2 *s2; +} str3; + +typedef struct { + str3 m; + str3 n; +} str4; + +int main(int argc, char* argv[]) +{ + const int N = 1024; + str4 p, *q; + int i; + + p.m.s2 = (str2 *) malloc (sizeof (str2)); + p.m.s2->f = (str1 *) malloc (sizeof (str1)); + p.m.s2->e = (int *) malloc (sizeof (int) * N); + p.m.s2->f->a = (int *) malloc (sizeof (int) * N); + p.m.s2->f->c = (int *) malloc (sizeof (int) * N); + + p.n.s2 = (str2 *) malloc (sizeof (str2)); + p.n.s2->f = (str1 *) malloc (sizeof (str1)); + p.n.s2->e = (int *) malloc (sizeof (int) * N); + p.n.s2->f->a = (int *) malloc (sizeof (int) * N); + p.n.s2->f->c = (int *) malloc (sizeof (int) * N); + + q = (str4 *) malloc (sizeof (str4)); + + q->m.s2 = (str2 *) malloc (sizeof (str2)); + q->m.s2->f = (str1 *) malloc (sizeof (str1)); + q->m.s2->e = (int *) malloc (sizeof (int) * N); + q->m.s2->f->a = (int *) malloc (sizeof (int) * N); + q->m.s2->f->c = (int *) malloc (sizeof (int) * N); + + q->n.s2 = (str2 *) malloc (sizeof (str2)); + q->n.s2->f = (str1 *) malloc (sizeof (str1)); + q->n.s2->e = (int *) malloc (sizeof (int) * N); + q->n.s2->f->a = (int *) malloc (sizeof (int) * N); + q->n.s2->f->c = (int *) malloc (sizeof (int) * N); + + for (i = 0; i < N; i++) + { + p.m.s2->e[i] = 0; + p.m.s2->f->a[i] = 0; + p.m.s2->f->c[i] = 0; + p.n.s2->e[i] = 0; + p.n.s2->f->a[i] = 0; + p.n.s2->f->c[i] = 0; + q->m.s2->e[i] = 0; + q->m.s2->f->a[i] = 0; + q->m.s2->f->c[i] = 0; + q->n.s2->e[i] = 0; + q->n.s2->f->a[i] = 0; + q->n.s2->f->c[i] = 0; + } + + for (int i = 0; i < 99; i++) + { + int j; +#pragma acc enter data copyin(p.m.s2[:1]) + { +#pragma acc parallel loop copy(p.m.s2->e[:N]) + for (j = 0; j < N; j++) + p.m.s2->e[j]++; + } +#pragma acc exit data delete(p.m.s2[:1]) + } + + for (i = 0; i < N; i++) + { + if (p.m.s2->e[i] != 99) + abort (); + } + + return 0; +} From patchwork Tue May 11 08:57:46 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1476987 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=2620:52:3:1:0:246e:9693:128c; helo=sourceware.org; envelope-from=gcc-patches-bounces@gcc.gnu.org; receiver=) Received: from 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 RSA-PSS (4096 bits) server-digest SHA256) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 4FfWzb6Mz1z9sSn for ; Tue, 11 May 2021 18:58:19 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id EAE6C388A413; Tue, 11 May 2021 08:58:05 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa2.mentor.iphmx.com (esa2.mentor.iphmx.com [68.232.141.98]) by sourceware.org (Postfix) with ESMTPS id AB2B4385780C; Tue, 11 May 2021 08:58:02 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org AB2B4385780C Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=Julian_Brown@mentor.com IronPort-SDR: oSm7LIou0skqstYomjQEhSq4aBw9RhVMfglTjTG0Y/zLPseYEwYX7tMyze/DCScMeXw+oJLnnH /A6BqyjdUn89RfrRlniLvR8dsr9rv4BefLfjmCfKwNgXJnCaTN2INuuwptf0gXnF/o1zloMOBx TrF0FEl0oyLpG5kcYO06SVM5ov2ttnJjVDthYa+qHCZvIh0+fLGr0RGgdwnvVRyTehqIUi8isp Gxodh9Ezn8bN8hDsHpD7r/KgSgmr2OmTXol/jLRs9I0l8vYaogHfjUJXUDiShp6G4OD7E+CCzY gqg= X-IronPort-AV: E=Sophos;i="5.82,290,1613462400"; d="scan'208";a="61053294" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa2.mentor.iphmx.com with ESMTP; 11 May 2021 00:58:02 -0800 IronPort-SDR: iE6p5ppkI3mej1Y4Kag+jyyflVxZMqFM1uPjW4qXc4LWLkTmW76KqAEUakfbza3voPblRfMXQP mDJ1PMMGb1d/zZ8MNf9w3yPt5nQny7WPglbAt1zofjeutppAi7K/lwgoUfwCrBiM6k6Jc/wGbI p4lrqHrgdo2WF5ocW9ZUI/derQJTkt77cTeS2TXSKgwq+cNLz9Ebes0+ZsmW6W9l7mfu+VP20r AOJ3gFxMHDwabeYDbFHHe5PoygJDKsUdclpXJDtW0wj3iwzThVC/OhqTHy1ipbIf7ISSzGItH/ GWA= From: Julian Brown To: Subject: [PATCH 7/7] [og10] WIP GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION changes Date: Tue, 11 May 2021 01:57:46 -0700 Message-ID: <04b90981e94acbabde004ee0992d404931e4fabf.1620721888.git.julian@codesourcery.com> X-Mailer: git-send-email 2.29.2 In-Reply-To: References: MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: SVR-IES-MBX-03.mgc.mentorg.com (139.181.222.3) To SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) X-Spam-Status: No, score=-12.0 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, RCVD_IN_DNSWL_NONE, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.2 X-Spam-Checker-Version: SpamAssassin 3.4.2 (2018-09-13) 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: , Cc: Jakub Jelinek , Tobias Burnus , fortran@gcc.gnu.org, Thomas Schwinge Errors-To: gcc-patches-bounces@gcc.gnu.org Sender: "Gcc-patches" This work-in-progress patch tries to get GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION to behave more like GOMP_MAP_ATTACH_DETACH -- in that the mapping is made to form groups to be processed by build_struct_group/build_struct_comp_map. I think that's important to integrate with how groups of mappings for array sections are handled in other cases. This patch isn't sufficient by itself to fix a couple of broken test cases at present (libgomp.c++/target-lambda-1.C, libgomp.c++/target-this-4.C), though. 2021-05-11 Julian Brown gcc/ * gimplify.c (build_struct_comp_nodes): Add GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION handling. (build_struct_group): Process GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION as part of pointer group. (gimplify_scan_omp_clauses): Update prev_list_p such that GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION will form part of pointer group. --- gcc/gimplify.c | 16 ++++++++++++---- 1 file changed, 12 insertions(+), 4 deletions(-) diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 6d204908c82..c5cb486aa23 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -8298,7 +8298,9 @@ build_struct_comp_nodes (enum tree_code code, tree grp_start, tree grp_end, if (grp_mid && OMP_CLAUSE_CODE (grp_mid) == OMP_CLAUSE_MAP && (OMP_CLAUSE_MAP_KIND (grp_mid) == GOMP_MAP_ALWAYS_POINTER - || OMP_CLAUSE_MAP_KIND (grp_mid) == GOMP_MAP_ATTACH_DETACH)) + || OMP_CLAUSE_MAP_KIND (grp_mid) == GOMP_MAP_ATTACH_DETACH + || (OMP_CLAUSE_MAP_KIND (grp_mid) + == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION))) { tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (grp_end), OMP_CLAUSE_MAP); @@ -8774,12 +8776,14 @@ build_struct_group (struct gimplify_omp_ctx *ctx, ? splay_tree_lookup (ctx->variables, (splay_tree_key) decl) : NULL); bool ptr = (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER); - bool attach_detach = (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH); + bool attach_detach = (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH + || (OMP_CLAUSE_MAP_KIND (c) + == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION)); bool attach = (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH); bool has_attachments = false; /* For OpenACC, pointers in structs should trigger an attach action. */ - if (attach_detach + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH && ((region_type & (ORT_ACC | ORT_TARGET | ORT_TARGET_DATA)) || code == OMP_TARGET_ENTER_DATA || code == OMP_TARGET_EXIT_DATA)) @@ -9784,6 +9788,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, if (!remove && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_POINTER && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH_DETACH + && (OMP_CLAUSE_MAP_KIND (c) + != GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION) && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET && OMP_CLAUSE_CHAIN (c) && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (c)) == OMP_CLAUSE_MAP @@ -9792,7 +9798,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c)) == GOMP_MAP_ATTACH_DETACH) || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c)) - == GOMP_MAP_TO_PSET))) + == GOMP_MAP_TO_PSET) + || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c)) + == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION))) prev_list_p = list_p; break;