From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from esa4.mentor.iphmx.com (esa4.mentor.iphmx.com [68.232.137.252]) by sourceware.org (Postfix) with ESMTPS id 0D7F63858D39; Wed, 6 Dec 2023 11:36:41 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 0D7F63858D39 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com ARC-Filter: OpenARC Filter v1.0.0 sourceware.org 0D7F63858D39 Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=68.232.137.252 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1701862605; cv=none; b=M+MLLTdQ97NCtY8qOjiaGocHKW554bHmAaNMsaelLLDd0yEOpB4+c4+bEeTJGnMpaaVNp8iuvF0agWaJJyxap5h2G6VcDQUjM9yiyIYnNt4cC0t1wpsIHdNNbBRUt8Ft5AmRJIx6RamyLrWtNjmKAYFGpDh6qXIJBe2UQvf+f6A= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1701862605; c=relaxed/simple; bh=VHSV5A/83WsJ4tyTc7v3j6wKtnOpTAza6ymWzw4LrIk=; h=Message-ID:Date:MIME-Version:Subject:To:From; b=iIl4DcmI/2z5OHv8f5whomcKGYGPn2fnIqVlMqem9Urwt0oDatM62eTq9LplUjNryLElTFjoLNT8c0KGezD3nONPa2N+QwPVpxM36NIZhysu4v4kIM6VojNY0oeIg9HD2iPOhYoFRFQ7sg/BOKL4hrlwAw0QLToqHZ4qkuR6Xjs= ARC-Authentication-Results: i=1; server2.sourceware.org X-CSE-ConnectionGUID: YMeEuQntRZq6gftelZQiQg== X-CSE-MsgGUID: 9t2UPkbaSN2WOU5msM3e1g== X-IronPort-AV: E=Sophos;i="6.04,255,1695715200"; d="scan'208";a="24598223" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa4.mentor.iphmx.com with ESMTP; 06 Dec 2023 03:36:41 -0800 IronPort-SDR: uXfVC0lK9b00GkTYFoip6qsysuF8Zb2wf0NNPMrurY1PSudY9tfwnrP47cc6FumwmAi7OjjQOt aRRaTsGpxl0DjfGh1fWqkob6TAkww235YsCJ5vJMOj65EWFOdJwqHh03QceRlnuZIA5AWVlLfd ivo6XVtHP6jo5JBx4BMTk4EOaa/eFRcHpBYD9mJNAOGRgTtr6jC5noul5FDzSyZTit+uICQcxp HZFXEGjfvWO4uucTWDAYUcGcIdilITh0CEjOdBJikrEwKVKfeBB8WlrzoUS9yR+T14aq3eYl8g BWE= Message-ID: Date: Wed, 6 Dec 2023 12:36:34 +0100 MIME-Version: 1.0 User-Agent: Mozilla Thunderbird Subject: Re: [PATCH v7 3/5] OpenMP: Pointers and member mappings Content-Language: en-US To: Julian Brown , CC: , References: From: Tobias Burnus In-Reply-To: Content-Type: text/plain; charset="UTF-8"; format=flowed Content-Transfer-Encoding: quoted-printable X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-10.mgc.mentorg.com (139.181.222.10) To svr-ies-mbx-12.mgc.mentorg.com (139.181.222.12) X-Spam-Status: No, score=-5.3 required=5.0 tests=BAYES_00,HEADER_FROM_DIFFERENT_DOMAINS,KAM_DMARC_STATUS,KAM_SHORT,SPF_HELO_PASS,SPF_PASS,TXREP,T_SCC_BODY_TEXT_LINE autolearn=no autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org List-Id: Hi Julian, LGTM, except for: * The 'target exit data' handling - comments below - looks a bit fishy/inconsistent. I intent to have a closer look with more code context, but maybe you should have a look at it as well. BTW: Fortran deep-mapping is not yet on mainline. Are you aware of changes or in particular testcases on OG13 related to your patch series that should be included when upstreaming that auto-mapping of allocatable components patch? * * * On 19.08.23 00:47, Julian Brown wrote: > This patch changes the mapping node arrangement used for array components > of derived types in order to accommodate for changes made in the previous > patch, particularly the use of "GOMP_MAP_ATTACH_DETACH" for pointer-typed > derived-type members instead of "GOMP_MAP_ALWAYS_POINTER". > > We change the mapping nodes used for a derived-type mapping like this: > > type T > integer, pointer, dimension(:) :: arrptr > end type T > > type(T) :: tvar > [...] > !$omp target map(tofrom: tvar%arrptr) > > So that the nodes used look like this: > > 1) map(to: tvar%arrptr) --> > GOMP_MAP_TO [implicit] *tvar%arrptr%data (the array data) > GOMP_MAP_TO_PSET tvar%arrptr (the descriptor) > GOMP_MAP_ATTACH_DETACH tvar%arrptr%data > > 2) map(tofrom: tvar%arrptr(3:8) --> > GOMP_MAP_TOFROM *tvar%arrptr%data(3) (size 8-3+1, etc.) [Clarification for the patch readers - not to the patch writer] At least to me the wording or implications of (1) were not fully clear to me. However, it is intended that the implicit mapping ensures that the pointer is actually mapped. Fortunately, that's indeed the case :-) as testing and also the new testcase libgomp/testsuite/libgomp.fortran/map-= subarray-4.f90 shows. The related part comes from OpenMP 5.1 [349:9-12] or slightly refined later, here the TR12 wording: If a list item in a map clause is an associated pointer that is not attach-= ineligible and the pointer is not the base pointer of another list item in a map clause on the same const= ruct, then it is treated as if its pointer target is implicitly mapped in the same clause. For the purpose= s of the map clause, the mapped pointer target is treated as if its base pointer is the associated p= ointer." [213:18-21] * * * > [...] > In the middle end, we process mappings using the struct sibling-list > handling machinery by moving the "GOMP_MAP_TO_PSET" node from the middle > of the group of three mapping nodes to the proper sorted position after > the GOMP_MAP_STRUCT mapping: > > GOMP_MAP_STRUCT tvar (len: 1) > GOMP_MAP_TO_PSET tvar%arr (size: 64, etc.) <--. moved here > [...] | > GOMP_MAP_TOFROM *tvar%arrptr%data(3) ___| > GOMP_MAP_ATTACH_DETACH tvar%arrptr%data > > In another case, if we have an array of derived-type values "dtarr", > and mappings like: > > i =3D 1 > j =3D 1 > map(to: dtarr(i)%arrptr) map(tofrom: dtarr(j)%arrptr(3:8)) > > We still map the same way, but this time we cannot prove that the base > expressions "dtarr(i) and "dtarr(j)" are the same in the front-end. > So we keep both mappings, but we move the "[implicit]" mapping of the > full-array reference to the end of the clause list in gimplify.cc (by > adjusting the topological sorting algorithm): > > GOMP_MAP_STRUCT dtvar (len: 2) > GOMP_MAP_TO_PSET dtvar(i)%arrptr > GOMP_MAP_TO_PSET dtvar(j)%arrptr > [...] > GOMP_MAP_TOFROM *dtvar(j)%arrptr%data(3) (size: 8-3+1) > GOMP_MAP_ATTACH_DETACH dtvar(j)%arrptr%data > GOMP_MAP_TO [implicit] *dtvar(i)%arrptr%data(1) (size: whole array) > GOMP_MAP_ATTACH_DETACH dtvar(i)%arrptr%data > > Always moving "[implicit]" full-array mappings after array-section > mappings (without that bit set) means that we'll avoid copying the whole > array unnecessarily -- even in cases where we can't prove that the arrays > are the same. > > The patch also fixes some bugs with "enter data" and "exit data" > directives with this new mapping arrangement. Also now if you have > mappings like this: > > #pragma omp target enter data map(to: dv, dv%arr(1:20)) > > The whole of the derived-type variable "dv" is mapped, so the > GOMP_MAP_TO_PSET for the array-section mapping can be dropped: > > GOMP_MAP_TO dv > > GOMP_MAP_TO *dv%arr%data > GOMP_MAP_TO_PSET dv%arr <-- deleted (array section mapping) > GOMP_MAP_ATTACH_DETACH dv%arr%data > > To accommodate for recent changes to mapping nodes made by > Tobias, this version of the patch avoids using GOMP_MAP_TO_PSET > for "exit data" directives, in favour of using the "correct" > GOMP_MAP_RELEASE/GOMP_MAP_DELETE kinds during early expansion. A new > flag is introduced so the middle-end knows when the latter two kinds > are being used specifically for an array descriptor. > > This version of the patch is based on the version posted for the og13 > branch: > > https://gcc.gnu.org/pipermail/gcc-patches/2023-June/622222.html [...] > + if (OMP_CLAUSE_MAP_KIND (node) =3D=3D GOMP_MAP_DELETE > + || OMP_CLAUSE_MAP_KIND (node) =3D=3D GOMP_MAP_RELEASE > + || op =3D=3D EXEC_OMP_TARGET_EXIT_DATA) > { > [...] > + gomp_map_kind map_kind > + =3D (op =3D=3D EXEC_OMP_TARGET_EXIT_DATA) ? GOMP_MAP_RELEASE > + : OMP_CLAUSE_MAP_KIND (node= ); > + OMP_CLAUSE_SET_MAP_KIND (node2, map_kind); > + OMP_CLAUSE_RELEASE_DESCRIPTOR (node2) =3D 1; For '!$omp target exit data map(delete: array)' this looks wrong as it replaces 'delete' by 'release' for the descriptor - while for '!$omp target (data) map(delete: array)' it remains 'delete'. Thus, I wonder whether that shouldn't be instead OMP_CLAUSE_MAP_KIND (node) =3D=3D GOMP_MAP_DELETE ? GOMP_MAP_DELETE : GOMP_MAP_RELEASE; * * * We later have the following; just reading the patch, I wonder whether GOMP_= TO_PSET is correct for a generic 'target exit data' here or not. It seems at a glance as if an "|| op =3D=3D 'target exit data'" is missing here: > - if (openacc) > - OMP_CLAUSE_SET_MAP_KIND (desc_node, > + if (openacc > + || (map_kind !=3D GOMP_MAP_RELEASE > + && map_kind !=3D GOMP_MAP_DELETE)) > + OMP_CLAUSE_SET_MAP_KIND (node2, > GOMP_MAP_TO_PSET); > else > - OMP_CLAUSE_SET_MAP_KIND (desc_node, map_kind); > - OMP_CLAUSE_DECL (desc_node) =3D inner; > - OMP_CLAUSE_SIZE (desc_node) =3D TYPE_SIZE_UNIT (t= ype); > - if (openacc) > - node2 =3D desc_node; > - else > + OMP_CLAUSE_SET_MAP_KIND (node2, map_kind); (Here, without setting OMP_CLAUSE_RELEASE_DESCRIPTOR.) And for completeness, we later have: > +omp_map_clause_descriptor_p (tree c) > +{ > + if (OMP_CLAUSE_CODE (c) !=3D OMP_CLAUSE_MAP) > + return false; > + > + if (OMP_CLAUSE_MAP_KIND (c) =3D=3D GOMP_MAP_TO_PSET) > + return true; > + > + if ((OMP_CLAUSE_MAP_KIND (c) =3D=3D GOMP_MAP_RELEASE > + || OMP_CLAUSE_MAP_KIND (c) =3D=3D GOMP_MAP_DELETE) > + && OMP_CLAUSE_RELEASE_DESCRIPTOR (c)) > + return true; * * * > + /* Save array descriptor for use > + in gfc_omp_deep_mapping{,_p,_cnt}; for= ce > + evaluate to ensure that it is > + not gimplified + is a decl. */ This is part of my Fortran allocatable-components deep-mapping patch that i= s currently only on OG9 (?) to OG13, badly needs to be upstreamed but required that Jak= ub had a look at it - well, I still would like that he has a look at the omp-low.cc = parts and it needs to be re-diffed. Hence, while I wouldn't mind to keep it to avoid more diffing work, I think= it will be cleaner not to keep it here. * * * Thanks, Tobias > 2023-08-18 Julian Brown > > gcc/fortran/ > * dependency.cc (gfc_omp_expr_prefix_same): New function. > * dependency.h (gfc_omp_expr_prefix_same): Add prototype. > * gfortran.h (gfc_omp_namelist): Add "duplicate_of" field to "u2= " > union. > * trans-openmp.cc (dependency.h): Include. > (gfc_trans_omp_array_section): Adjust mapping node arrangement f= or > array descriptors. Use GOMP_MAP_TO_PSET or > GOMP_MAP_RELEASE/GOMP_MAP_DELETE with the OMP_CLAUSE_RELEASE_DES= CRIPTOR > flag set. > (gfc_symbol_rooted_namelist): New function. > (gfc_trans_omp_clauses): Check subcomponent and subarray/element > accesses elsewhere in the clause list for pointers to derived ty= pes or > array descriptors, and adjust or drop mapping nodes appropriatel= y. > Adjust for changes to mapping node arrangement. > (gfc_trans_oacc_executable_directive): Pass code op through. > > gcc/ > * gimplify.cc (omp_map_clause_descriptor_p): New function. > (build_omp_struct_comp_nodes, omp_get_attachment, omp_group_base= ): Use > above function. > (omp_tsort_mapping_groups): Process nodes that have > OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P set after those that don't. A= dd > enter_exit_data parameter. > (omp_resolve_clause_dependencies): Remove GOMP_MAP_TO_PSET mappi= ngs if > we're mapping the whole containing derived-type variable. > (omp_accumulate_sibling_list): Adjust GOMP_MAP_TO_PSET handling. > Remove GOMP_MAP_ALWAYS_POINTER handling. > (gimplify_scan_omp_clauses): Pass enter_exit argument to > omp_tsort_mapping_groups. Don't adjust/remove GOMP_MAP_TO_PSET > mappings for derived-type components here. > * tree.h (OMP_CLAUSE_RELEASE_DESCRIPTOR): New macro. > > gcc/testsuite/ > * gfortran.dg/gomp/map-9.f90: Adjust scan output. > * gfortran.dg/gomp/map-subarray-2.f90: New test. > * gfortran.dg/gomp/map-subarray.f90: New test. > > libgomp/ > * testsuite/libgomp.fortran/map-subarray.f90: New test. > * testsuite/libgomp.fortran/map-subarray-2.f90: New test. > * testsuite/libgomp.fortran/map-subarray-3.f90: New test. > * testsuite/libgomp.fortran/map-subarray-4.f90: New test. > * testsuite/libgomp.fortran/map-subarray-6.f90: New test. > * testsuite/libgomp.fortran/map-subarray-7.f90: New test. > * testsuite/libgomp.fortran/map-subarray-8.f90: New test. > * testsuite/libgomp.fortran/map-subcomponents.f90: New test. > * testsuite/libgomp.fortran/struct-elem-map-1.f90: Adjust for > descriptor-mapping changes. Remove XFAIL. ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstra=C3=9Fe 201= , 80634 M=C3=BCnchen; Gesellschaft mit beschr=C3=A4nkter Haftung; Gesch=C3= =A4ftsf=C3=BChrer: Thomas Heurung, Frank Th=C3=BCrauf; Sitz der Gesellschaf= t: M=C3=BCnchen; Registergericht M=C3=BCnchen, HRB 106955