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 536CB3858005; Tue, 14 Nov 2023 10:21:38 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 536CB3858005 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 536CB3858005 Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=68.232.137.180 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1699957300; cv=none; b=V4yxfaafSja8eOOnFFEPrLtzIWW4PEPIRBLKq1miNwZpZv4T7mp3NIKhBR8LGedruJNEVQpZsixGxI1YXPDiHkHRHAloKldmGmzPP3q+vUxKb3ycHzoV1E/kuiORH+2EnVZZ44v6EvEhjdOY3q4F+7Cr9iSWObRMmh9gINfGoFE= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1699957300; c=relaxed/simple; bh=2+JpneSwQXDi5tkBlnZJnC6/+NSEceokzpMsIojq8s8=; h=Message-ID:Date:MIME-Version:Subject:To:From; b=o8TKIDK9O7Z0nIVrtNVarMNR44hxz1dUSLA+7It7j0FVZ+1n4Tq2XwsH6JSwo9YV8GX2yD3Zr2WHNKQVEW/0skSL/N+ZX4ZUhQApBySJ88I4C+mVipFQZXCZuqAzx0zBP4vkQTZZ7S2VUdj2+K+mGubc0ldVVw1ve/dhT0kCIaU= ARC-Authentication-Results: i=1; server2.sourceware.org X-CSE-ConnectionGUID: oe2BiJ8aQZOPfG4zGW+o9Q== X-CSE-MsgGUID: IpiPtULfTo6woAlkYMw5ow== X-IronPort-AV: E=Sophos;i="6.03,301,1694764800"; d="scan'208";a="22717192" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa3.mentor.iphmx.com with ESMTP; 14 Nov 2023 02:21:36 -0800 IronPort-SDR: aVXiVbbKNutraxMDekVTAAtUkibWHKbLcWJYERk7y2ntBod4KUkAoK55Qus3o8K3yeSuoFHCc4 G2YyNyNOy5CH488i1NbmGwHy/IDjY2PS8YEsstpHF+mMJ710fVypH+Yt7Etja08kczUQzC2bai 1v5ZTJysuYe1IOIdyYOd85PDSJIMS2GVw7prpUf8lESCMDNB9Eb5jFnJPFRbF1WOe9T9pF51q6 Hv82kyLjK1JuhIVVqISIZwyc8JOnz/inL9d7kotBIqR7Mskg5as9JtBihGZlg387WD67fDLw1h hg0= Message-ID: Date: Tue, 14 Nov 2023 11:21:30 +0100 MIME-Version: 1.0 User-Agent: Mozilla Thunderbird Subject: Re: [PATCH v7 2/5] OpenMP/OpenACC: Rework clause expansion and nested struct handling 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-14.mgc.mentorg.com (139.181.222.14) 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, first round of comments - I think I need a second pass as the patch is long and complex. The depth of review also was decreasing, hence, I assume I will spot things in later parts of the compiler. In any case, I think the patch is a huge leap forward and very useful! Contrary to previous iterations of the patch series, it looks as if {1,2}/5 could be sensibly applied before {3,4,5}/5. However, those are smaller or tiny, it should be way quicker to review them! * * * CROSS REFS This patch is part of the patch set "[PATCH v7 0/5] OpenMP/OpenACC: map clause and OMP gimplify rework" https://gcc.gnu.org/pipermail/gcc-patches/2023-August/627895.html This patch set in turn is required for the follow-up patch set https://gcc.gnu.org/pipermail/gcc-patches/2023-September/629363.html '[PATCH 0/8] OpenMP: lvalue parsing and "declare mapper" support' and it looks as if it would be very useful for some other WIP patches like map+iterator. * * * The patch 1/5 is just an indentation patch to make the following patch smaller - and it has been approved multiple times. This patch can be found at https://gcc.gnu.org/pipermail/gcc-patches/2023-August/627897.html in case someone wants to read the full patch or Julian's patch-intro comments. * * * I do note that all 5 patches of this patch set apply cleanly with some fuzzy line matching, but: WARNING: at least for this patch (2/5), code gets wrongly applied in (c-)parser.cc - see below. On 19.08.23 00:47, Julian Brown wrote: > This patch reworks clause expansion in the C, C++ and (to a lesser > extent) Fortran front ends for OpenMP and OpenACC mapping nodes used in > GPU offloading support. > > At present a single clause may be turned into several mapping nodes, > or have its mapping type changed, in several places scattered through > the front- and middle-end. The analysis relating to which particular > transformations are needed for some given expression has become quite har= d > to follow. Briefly, we manipulate clause types in the following places: [...] [As the patch is very large, I copied the interesting bits out of the patch and comment on them - I hope that preserves enough context to be useful.] tree c_omp_address_inspector::get_address () This class member function is unused. It returns 'orig' and that member var= iable is used a couple of times in the member functions - but not this function. I think it can be removed. (I also find the name slightly misleading.) * * * +c_omp_address_inspector::maybe_zero_length_array_section (tree clause) +{ + switch (OMP_CLAUSE_MAP_KIND (clause)) + { + case GOMP_MAP_ALLOC: + case GOMP_MAP_IF_PRESENT: + case GOMP_MAP_TO: + case GOMP_MAP_FROM: + case GOMP_MAP_TOFROM: + case GOMP_MAP_ALWAYS_TO: + case GOMP_MAP_ALWAYS_FROM: + case GOMP_MAP_ALWAYS_TOFROM: + case GOMP_MAP_RELEASE: + case GOMP_MAP_DELETE: + case GOMP_MAP_FORCE_TO: + case GOMP_MAP_FORCE_FROM: + case GOMP_MAP_FORCE_TOFROM: + case GOMP_MAP_FORCE_PRESENT: Shouldn't this also include: GOMP_MAP_ALWAYS_PRESENT_TO: GOMP_MAP_ALWAYS_PRESENT_FROM: GOMP_MAP_ALWAYS_PRESENT_TOFROM: GOMP_MAP_PRESENT_ALLOC: GOMP_MAP_PRESENT_TO: ? * * * +omp_expand_access_chain (tree c, tree expr, vec &addr_to= kens, + unsigned *idx) +{ + using namespace omp_addr_tokenizer; + location_t loc =3D OMP_CLAUSE_LOCATION (c); + unsigned i =3D *idx; + tree c2 =3D NULL_TREE; + gomp_map_kind kind + =3D (OMP_CLAUSE_CODE (c) =3D=3D OMP_CLAUSE_FROM + || (OMP_CLAUSE_CODE (c) =3D=3D OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (c) =3D=3D GOMP_MAP_FROM + || OMP_CLAUSE_MAP_KIND (c) =3D=3D GOMP_MAP_DELETE + || OMP_CLAUSE_MAP_KIND (c) =3D=3D GOMP_MAP_RELEASE + || OMP_CLAUSE_MAP_KIND (c) =3D=3D GOMP_MAP_ALWAYS_FROM + || OMP_CLAUSE_MAP_KIND (c) =3D=3D GOMP_MAP_FORCE_FROM + || OMP_CLAUSE_MAP_KIND (c) =3D=3D GOMP_MAP_PRESENT_FROM))) + ? GOMP_MAP_DETACH : GOMP_MAP_ATTACH; GOMP_MAP_ALWAYS_PRESENT_FROM missing? And I personally find it more readable to have if (... ... ...) kind =3D GOMP_MAP_DETACH; else kind =3D GOMP_MAP_ATTACH; as in the assignment above, the lhs is in this case really far from the rhs= . * * * I do see in omp_expand_access_chain, which is shared by C and C++, that the= following accesses are handled: ACCESS_POINTER, ACCESS_POINTER_OFFSET, ACCESS_INDEXED_ARRAY, But not: ACCESS_REF_TO_POINTER, ACCESS_REF_TO_POINTER_OFFSET, ACCESS_INDEXED_REF_TO_ARRAY Do those need to be handled as well? SIDENOTE: As Fortran allocatable components, reference-type variables in a = C++ class ('class', 'struct') are automatically be copied when doing 'map(var)' - contrary to pointer mem= bers that need to be explicitly be mapped (either at mapping side or via a mapper). (And contrary to Fortra= n allocatables, reference-type members are always "there" - not like with allocatables can change from unallocated= to allocated - with all compilations of polymorphism, array sizes, string lengths, ...). * * * +tree +c_omp_address_inspector::expand_array_base (tree c, ... + bool map_p =3D OMP_CLAUSE_CODE (c) =3D=3D OMP_CLAUSE_MAP; + bool implicit_p =3D (OMP_CLAUSE_CODE (c) =3D=3D OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_IMPLICIT (c)); ... + if (!openmp + && OMP_CLAUSE_CODE (c) =3D=3D OMP_CLAUSE_MAP How about using 'map_p' (twice), given that it is already defined? ... + if (target) + { + c2 =3D build_omp_clause (loc, OMP_CLAUSE_MAP); + if (target + && !ref_to_ptr Remove 'target &&' from the second condition. * * * +++ b/gcc/c/c-parser.cc @@ -17659,7 +17659,8 @@ c_parser_omp_clause_detach (c_parser *parser, tree = list) static tree c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, - const char *where, bool finish_p =3D true) + const char *where, bool finish_p =3D true, + bool target =3D false) Here but also in gcc/c-family/c-omp.cc - I wonder whether this should be ta= rget_p for consistency with most other boolean variables. The code is not really consi= stent, but given that there is 'finish_p' here and map_p etc. in the other file, i= t seems to make sense to use 'target_p. * * * If I try the example at https://discourse.llvm.org/t/how-do-i-map-an-array-of-pointers-using-arr= ay-notation/53640/3 it fails to compile with: foo.c:11:34: error: =E2=80=98#pragma omp target enter data=E2=80=99 with ma= p-type other than =E2=80=98to=E2=80=99, =E2=80=98tofrom=E2=80=99 or =E2=80= =98alloc=E2=80=99 on =E2=80=98map=E2=80=99 clause 11 | #pragma omp target enter data map(to: a[i][0:1]) The problem is that fuzzy 'patch' mis-apply the patch for c-parser.cc's @@ -22106,6 +22111,7 @@ c_parser_omp_target_enter_data (location_t loc, c_p= arser *parser, case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_ALWAYS_POINTER: case GOMP_MAP_ATTACH_DETACH: + case GOMP_MAP_ATTACH: break; default: map_seen |=3D 1; @@ -22215,6 +22221,7 @@ c_parser_omp_target_exit_data (location_t loc, c_pa= rser *parser, case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_ALWAYS_POINTER: case GOMP_MAP_ATTACH_DETACH: + case GOMP_MAP_DETACH: break; default: map_seen |=3D 1; Namely after applying it, the result is: @@ -23329,0 +23335 @@ c_parser_omp_target_data (location_t loc, c_parser *p= arser, bool *if_p) + case GOMP_MAP_ATTACH: @@ -23494,0 +23501 @@ c_parser_omp_target_enter_data (location_t loc, c_par= ser *parser, + case GOMP_MAP_DETACH: That is, the new code is added to c_parser_omp_target_data and c_parser_omp_target_enter_data instead of to c_parser_omp_target_enter_data and c_parser_omp_target_exit_data. * * * Likewise issue for C++'s cp/parser.cc. * * * The first question is whether such code is also needed for 'target' and 'ta= rget data' besides 'target enter/exit data'. It looks as if omp_expand_access_chain sh= ould also apply here. * * * And trying the following snippet: void f(){ int **a,i,j,n; #pragma omp target enter data map(a[i][j:n]) #pragma omp target exit data map(a[i][j:n]) #pragma omp target data map(a[i][j:n]) { #pragma omp target map(a[i][j:n]) ; } } Results - when putting the enums to enter/exit data instead of to target data / target enter data: g++ fails, unless I messed up, with: fiof.c:4:36: error: =E2=80=98#pragma omp target exit data=E2=80=99 with map= -type other than =E2=80=98from=E2=80=99, =E2=80=98tofrom=E2=80=99, =E2=80= =98release=E2=80=99 or =E2=80=98delete=E2=80=99 on =E2=80=98map=E2=80=99 cl= ause 4 | #pragma omp target exit data map(a[i][j:n]) | ^ fiof.c:5:31: error: =E2=80=98#pragma omp target data=E2=80=99 with map-type= other than =E2=80=98to=E2=80=99, =E2=80=98from=E2=80=99, =E2=80=98tofrom= =E2=80=99 or =E2=80=98alloc=E2=80=99 on =E2=80=98map=E2=80=99 clause 5 | #pragma omp target data map(a[i][j:n]) And gcc fails with: fiof.c:3:36: error: =E2=80=98a=E2=80=99 appears more than once in map claus= es 3 | #pragma omp target enter data map(a[i][j:n]) | ^ fiof.c:4:35: error: =E2=80=98a=E2=80=99 appears more than once in map claus= es 4 | #pragma omp target exit data map(a[i][j:n]) | ^ fiof.c:5:30: error: =E2=80=98a=E2=80=99 appears more than once in map claus= es 5 | #pragma omp target data map(a[i][j:n]) | ^ fiof.c:7:25: error: =E2=80=98a=E2=80=99 appears more than once in map claus= es 7 | #pragma omp target map(a[i][j:n]) (Plus all '... data' ones: '... must contain at least one ...') Can you fix this - and add an example in the spirit of the LLVM example abo= ve, except that it should test all four 'target' variants. Thanks, Tobias ----------------- 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