From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from esa3.mentor.iphmx.com (esa3.mentor.iphmx.com [68.232.137.180]) by sourceware.org (Postfix) with ESMTPS id 3CF9C3858C54; Sun, 18 Sep 2022 19:10:18 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 3CF9C3858C54 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,325,1654588800"; d="scan'208,223";a="83054834" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa3.mentor.iphmx.com with ESMTP; 18 Sep 2022 11:10:17 -0800 IronPort-SDR: Eu5nIVExAWzHc00o2ycEAOTcP9JsJCSSeAnwo4j2ZDKRaKVArvDybbiiplli1Qe8pU+vyOU0bM irFSP5+Jjs248FVi8zzQy+dvA+0zUpI8GoMW4Xgw4Qo49od6KwclCIQb+gsAiX+vZBJldQOnfu pPTIcv73JqiwnxQt+UBcnIifOgdoKMxHSJy5uKxTjluOq7wJ8YlsxevUF8tfrrWMACdC70IaV/ p5E2BnDJ8J8fIosQt0onz/xejQYZueC2GHltqIf8+qXSbtLtI8zCn+1M18kn4ukvlX7KMr2P+8 lsw= Date: Sun, 18 Sep 2022 20:10:00 +0100 From: Julian Brown To: Jakub Jelinek CC: , , , Subject: Re: [PATCH v3 05/11] OpenMP: push attaches to end of clause list in "target" regions Message-ID: <20220918201000.0138193c@squid.athome> In-Reply-To: References: <479bff9d51ee4db1ff46e0edaaf24d2a601f7a0d.1663101299.git.julian@codesourcery.com> Organization: Mentor Graphics X-Mailer: Claws Mail 4.0.0 (GTK+ 3.24.31; x86_64-pc-linux-gnu) MIME-Version: 1.0 Content-Type: multipart/mixed; boundary="MP_/a2kZyWhfbj9YoA8am5ndAUE" X-ClientProxiedBy: svr-orw-mbx-12.mgc.mentorg.com (147.34.90.212) To svr-ies-mbx-11.mgc.mentorg.com (139.181.222.11) X-Spam-Status: No, score=-10.4 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.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org List-Id: --MP_/a2kZyWhfbj9YoA8am5ndAUE Content-Type: text/plain; charset="US-ASCII" Content-Transfer-Encoding: 7bit Content-Disposition: inline On Wed, 14 Sep 2022 14:44:54 +0200 Jakub Jelinek wrote: > On Tue, Sep 13, 2022 at 02:03:15PM -0700, Julian Brown wrote: > > 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. > > Shouldn't this be done at the end of gimplify_adjust_omp_clauses? > I mean, can't further attach clauses appear because of declare mapper > for implicitly mapped variables? > Other than that, it is yet another walk of the whole clause list, so > would be nicer if it could be done in an existing walk over the > clauses or at least have a flag whether there are any such clauses > present and do it only in that case. > If it could be done in the main gimplify_adjust_omp_clauses loop, > nice, if it can appear also during > splay_tree_foreach (ctx->variables, gimplify_adjust_omp_clauses_1, > &data); that isn't the case. I don't think any ATTACH clauses can appear during the gimplify_adjust_omp_clause_1 walk. So, how about this? Thanks, Julian --MP_/a2kZyWhfbj9YoA8am5ndAUE Content-Type: text/x-patch Content-Transfer-Encoding: 7bit Content-Disposition: attachment; filename="0001-OpenMP-Push-attaches-to-end-of-clause-list-in-target.patch" >From d583f3315ac8cda58bbfce8c8574d0adc5283b00 Mon Sep 17 00:00:00 2001 From: Julian Brown Date: Wed, 7 Sep 2022 21:45:07 +0000 Subject: [PATCH 1/2] OpenMP: Push attaches to end of clause list in "target" regions 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. This version of the patch moves the attach-node motion to gimplify_adjust_omp_clauses. 2022-09-15 Julian Brown gcc/ * gimplify.cc (omp_segregate_mapping_groups): Update comment. (gimplify_adjust_omp_clauses): Move ATTACH and ATTACH_ZERO_LENGTH_ARRAY_SECTION nodes to the end of the clause list for offloaded OpenMP regions. gcc/testsuite/ * g++.dg/gomp/target-lambda-1.C: Adjust expected scan output. --- gcc/gimplify.cc | 37 ++++++++++++++++++++- gcc/testsuite/g++.dg/gomp/target-lambda-1.C | 2 +- 2 files changed, 37 insertions(+), 2 deletions(-) diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index 2ae0c8cb250..4d032c6bf06 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -9639,7 +9639,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. + Note that ATTACH nodes are later moved to the end of the list in + gimplify_adjust_omp_clauses, for target regions. */ static omp_mapping_group * omp_segregate_mapping_groups (omp_mapping_group *inlist) @@ -12411,10 +12413,15 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, *list_p = c2; } } + + tree attach_list = NULL_TREE; + tree *attach_tail = &attach_list; + while ((c = *list_p) != NULL) { splay_tree_node n; bool remove = false; + bool move_attach = false; switch (OMP_CLAUSE_CODE (c)) { @@ -12576,6 +12583,19 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, remove = true; break; } + /* 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). + 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). */ + if ((ctx->region_type & ORT_TARGET) != 0 + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH + || (OMP_CLAUSE_MAP_KIND (c) + == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION))) + move_attach = true; decl = OMP_CLAUSE_DECL (c); /* Data clauses associated with reductions must be compatible with present_or_copy. Warn and adjust the clause @@ -12890,10 +12910,25 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, if (remove) *list_p = OMP_CLAUSE_CHAIN (c); + else if (move_attach) + { + /* Remove attach node from here, separate out into its own list. */ + *attach_tail = c; + *list_p = OMP_CLAUSE_CHAIN (c); + OMP_CLAUSE_CHAIN (c) = NULL_TREE; + attach_tail = &OMP_CLAUSE_CHAIN (c); + } else list_p = &OMP_CLAUSE_CHAIN (c); } + /* Splice attach nodes at the end of the list. */ + if (attach_list) + { + *list_p = attach_list; + list_p = attach_tail; + } + /* Add in any implicit data sharing. */ struct gimplify_adjust_omp_clauses_data data; if ((gimplify_omp_ctxp->region_type & ORT_ACC) == 0) 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" } } */ -- 2.29.2 --MP_/a2kZyWhfbj9YoA8am5ndAUE--