From patchwork Tue Sep 13 21:03:15 2022 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1194 Return-Path: Delivered-To: ouuuleilei@gmail.com Received: by 2002:a5d:5044:0:0:0:0:0 with SMTP id h4csp2530244wrt; Tue, 13 Sep 2022 14:08:00 -0700 (PDT) X-Google-Smtp-Source: AA6agR4L1sU8uOno2kIRR6VVm4qI0Fb6KuWDCJg7c3P17SOjd1TQ1iMWwiAO1i8O+4mJs9Iqc9u7 X-Received: by 2002:a17:907:6ea3:b0:77c:4e23:9b2d with SMTP id sh35-20020a1709076ea300b0077c4e239b2dmr10115270ejc.6.1663103280435; Tue, 13 Sep 2022 14:08:00 -0700 (PDT) ARC-Seal: i=1; a=rsa-sha256; t=1663103280; cv=none; d=google.com; s=arc-20160816; b=RLchX8bkC088oOJpNJgaIsbb4ylhvn5hUXfQjBxivuQVlxRpJgD6vuF/ZS8hJCD46s T8xkwSQjp0wVwy2h082Icf0SinrFtCwArGp+wE05fy1Ye1O9oMfV1jOeMAr+lmK1RwLo HXBTv1oEiKxy//QdYxVH/Wx3OaBJ8U7YqWF35fDCGLFi0afHbya/ayXngbX482cObHpv FbCE28hEZfPivLu1zlekymLkF4HlMU2uQtee6m7htq1KxIUvlJ8ZATpvY9UCTb4agzG1 mz49te6/wzPp0CNDyMHif62aFKXCYdgySniQiHw3bGD4iyYcNuHFoQy0BR1iMFqFQedn iItA== ARC-Message-Signature: i=1; a=rsa-sha256; c=relaxed/relaxed; d=google.com; s=arc-20160816; h=sender:errors-to:cc:list-subscribe:list-help:list-post:list-archive :list-unsubscribe:list-id:precedence:content-transfer-encoding :mime-version:references:in-reply-to:message-id:date:subject:to:from :ironport-sdr:dmarc-filter:delivered-to; bh=pqJYw17d199Z/jzdDLx5AOQnemzcjIaDivH5VlkbesU=; b=JWrJ+4Amobq7MLh/GOtnbHT9LrOwa7qOMgeVLZo2RHlJMIMHSSOJdPu+BNRzogklx6 pETBTK5zbvvgj0itspy9Si47Ib6Lpn8isv83ZVKPr9qI3xGVRD03sS3Vy7y1xc4iUDKT kov7vvEeyMQysVtb4qOr48JaP2aSbYjf+luC/TCV6fuvYQUkpa9s4IVeoOAEOWB1Z/Qj 22Es93lQgfoT8bF3ERVWNUJCK+e0ugvY1nRo7SiksXKhLVkMrvePgyELFXQbmPAPdf6d opYOTlOed3DXTv6IPVjrZgyJMq9kPvDUQoPr2VIuPEbQk/tqIjk7kCsjloL6omGnc5fL HYwg== ARC-Authentication-Results: i=1; mx.google.com; spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 8.43.85.97 as permitted sender) smtp.mailfrom="gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org" Received: from sourceware.org (ip-8-43-85-97.sourceware.org. [8.43.85.97]) by mx.google.com with ESMTPS id t29-20020a50ab5d000000b0044ebbd162bfsi9128592edc.283.2022.09.13.14.08.00 for (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Tue, 13 Sep 2022 14:08:00 -0700 (PDT) Received-SPF: pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 8.43.85.97 as permitted sender) client-ip=8.43.85.97; Authentication-Results: mx.google.com; spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 8.43.85.97 as permitted sender) smtp.mailfrom="gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org" Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 9B7E73954434 for ; Tue, 13 Sep 2022 21:04:07 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa4.mentor.iphmx.com (esa4.mentor.iphmx.com [68.232.137.252]) by sourceware.org (Postfix) with ESMTPS id 12A30383A22C; Tue, 13 Sep 2022 21:03:30 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 12A30383A22C Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com X-IronPort-AV: E=Sophos;i="5.93,313,1654588800"; d="scan'208";a="82933007" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa4.mentor.iphmx.com with ESMTP; 13 Sep 2022 13:03:31 -0800 IronPort-SDR: 21+13hJszIu3IpD53urbvgtTIHkFbBioziHA+ePdPkVrTRq+AJf6pn+YtUymfLeRRZ8TWWqFPo xA4kJJI3usSGEUlqw3Dyxy3Y2TD6x5HdXFUod1ZBG7Tx4NAx3XSpfn4XUU4YKdF0UwOlnbKm/0 5KVKJg4J6DbsOfNnkEWTEQU10TtZvNTlXmO1N+6NE6pUOHxPlMNgLdwxXuFkjsTPfPLZkeIFN5 Y5fbjtJEzhJn/K3ArZ3ozOMDAtLgBViP5j994dRbgkPNgRT8QplFo5JF2qULN/4hyfeq/lRhEi 83E= From: Julian Brown To: Subject: [PATCH v3 05/11] OpenMP: push attaches to end of clause list in "target" regions Date: Tue, 13 Sep 2022 14:03:15 -0700 Message-ID: <479bff9d51ee4db1ff46e0edaaf24d2a601f7a0d.1663101299.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-13.mgc.mentorg.com (139.181.222.13) To svr-ies-mbx-11.mgc.mentorg.com (139.181.222.11) X-Spam-Status: No, score=-11.7 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP, T_SCC_BODY_TEXT_LINE autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Cc: Jakub Jelinek , tobias@codesourcery.com, fortran@gcc.gnu.org Errors-To: gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org Sender: "Gcc-patches" X-getmail-retrieved-from-mailbox: =?utf-8?q?INBOX?= X-GMAIL-THRID: =?utf-8?q?1743890185379682165?= X-GMAIL-MSGID: =?utf-8?q?1743890185379682165?= This patch moves GOMP_MAP_ATTACH{_ZERO_LENGTH_ARRAY_SECTION} nodes to the end of the clause list, for offload regions. This ensures that when we do the attach operation, both the "attachment point" and the target region have both already been mapped on the target. This avoids a pathological case that can otherwise happen with struct sibling-list handling. 2022-09-13 Julian Brown gcc/ * gimplify.cc (omp_segregate_mapping_groups): Update comment. (omp_push_attaches_to_end): New function. (gimplify_scan_omp_clauses): Use omp_push_attaches_to_end for offloaded regions. --- gcc/gimplify.cc | 66 ++++++++++++++++++++- gcc/testsuite/g++.dg/gomp/target-lambda-1.C | 2 +- 2 files changed, 65 insertions(+), 3 deletions(-) diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index c7998c2ccbd..bc7848843b3 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -9641,7 +9641,9 @@ omp_tsort_mapping_groups (vec *groups, /* Split INLIST into two parts, moving groups corresponding to ALLOC/RELEASE/DELETE mappings to one list, and other mappings to another. The former list is then appended to the latter. Each sub-list retains the - order of the original list. */ + order of the original list. + See also omp_push_attaches_to_end below -- we call that later after scanning + omp clauses. */ static omp_mapping_group * omp_segregate_mapping_groups (omp_mapping_group *inlist) @@ -9681,6 +9683,55 @@ omp_segregate_mapping_groups (omp_mapping_group *inlist) return tf_groups; } +/* This function moves GOMP_MAP_ATTACH{_ZERO_LENGTH_ARRAY_SECTION} nodes to the + end of the clause list, for offload regions. This ensures that when we do + the attach, both the "attachment point" and the target region have both + already been mapped on the target. This avoids a pathological case that can + otherwise happen with struct sibling-list handling. + + Do not call this for non-offload regions, e.g. for "enter data" or + "exit data" directives. + + The order of attach nodes and of non-attach nodes is otherwise retained. */ + +static tree +omp_push_attaches_to_end (tree list) +{ + tree nonattach_list = NULL_TREE, attach_list = NULL_TREE; + tree *nonattach_tail = &nonattach_list, *attach_tail = &attach_list; + + for (tree w = list; w;) + { + tree next = OMP_CLAUSE_CHAIN (w); + + if (OMP_CLAUSE_CODE (w) != OMP_CLAUSE_MAP) + goto nonattach; + + switch (OMP_CLAUSE_MAP_KIND (w)) + { + case GOMP_MAP_ATTACH: + case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION: + *attach_tail = w; + OMP_CLAUSE_CHAIN (w) = NULL_TREE; + attach_tail = &OMP_CLAUSE_CHAIN (w); + break; + + default: + nonattach: + *nonattach_tail = w; + OMP_CLAUSE_CHAIN (w) = NULL_TREE; + nonattach_tail = &OMP_CLAUSE_CHAIN (w); + } + + w = next; + } + + /* Splice lists together. */ + *nonattach_tail = attach_list; + + return nonattach_list; +} + /* Given a list LIST_P containing groups of mappings given by GROUPS, reorder those groups based on the output list of omp_tsort_mapping_groups -- singly-linked, threaded through each element's NEXT pointer starting at @@ -11950,7 +12001,18 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, list_p = &OMP_CLAUSE_CHAIN (c); } - ctx->clauses = *orig_list_p; + if ((region_type & ORT_TARGET) != 0) + /* If we have a target region, we can push all the attaches to the end of + the list (we may have standalone "attach" operations synthesized for + GOMP_MAP_STRUCT nodes that must be processed after the attachment point + AND the pointed-to block have been mapped). */ + ctx->clauses = omp_push_attaches_to_end (*orig_list_p); + else + /* ...but if we have something else, e.g. "enter data", we need to keep + "attach" nodes together with the previous node they attach to so that + separate "exit data" operations work properly (see libgomp/target.c). */ + ctx->clauses = *orig_list_p; + gimplify_omp_ctxp = ctx; } diff --git a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C index bff7fa7c669..5ce8ceadb19 100644 --- a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C +++ b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C @@ -87,7 +87,7 @@ int main (void) return 0; } -/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(b\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:iptr \[pointer assign, bias: 0\]\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) 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\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:__closure->__iptr \[bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\)} "gimple" } } */ +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(b\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:iptr \[pointer assign, bias: 0\]\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) 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\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:__closure->__iptr \[bias: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\)} "gimple" } } */ /* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(end\) firstprivate\(begin\) map\(to:loop \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:loop\.__data1 \[bias: 0\]\)} "gimple" } } */