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 E5C27386185D for ; Thu, 15 Jul 2021 10:33:16 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org E5C27386185D Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com IronPort-SDR: 0cJ6huN/74DdQMW4qmxGU7xKQessyWB4t83oYZolJ7BW7Y0g8KpPwB+gWeQtEYZUGnxVdVFgfp D2DnwZ/YO2HEgqVT3cwzZVybbwWrdTjaPCfMIzje+p3K0GK3be9x+QKBPXRwRlu5MUbhld4mdA oqzdkhRZval+pl7R34meOP8HT1CT4uhkyfPgX89am/gfSiiuVXkYrIickS1Gb9nM6bc/E6nSeV ZhdjwLO4zfp98nhwtGH4Y+honUshB0xcTf3SLZHF9yBh6W2jHcTqWQc6dCnXGtHf2Lx1IYmAQm 07A= X-IronPort-AV: E=Sophos;i="5.84,240,1620720000"; d="scan'208";a="63718507" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa4.mentor.iphmx.com with ESMTP; 15 Jul 2021 02:33:13 -0800 IronPort-SDR: u1PsQJG4fh2bWwjYvaxCssHsOkxwCu1ZNa1QzLmEpvbJVIZISX60no2416D7HTt3t7UUhmHbu7 qZoYUglrjp/FCxRrLMLvY5KukSWPE2p+ZVDGVtcS26L4rfbtmfbfJxQglv8wtkUtHfg3roPrSd un2VBdqvbAIexDidu9lk9pDqOuSw8ntNIyhHCqjmMppC2i2EKCUH924objP6Jjhl8u0+ZPOLSg nw6cl5JodtbcPv2UgcRw3InPn18atVEBju7blq7699+cGzEq9HrG1pSBFx/owWdEXNww2t2oBY gXI= From: Thomas Schwinge To: Richard Biener , Abid Qadeer , Jakub Jelinek CC: GCC Patches Subject: Re: [PATCH] [DWARF] Fix hierarchy of debug information for offload kernels. In-Reply-To: References: <20210701151657.935006-1-abidh@codesourcery.com> User-Agent: Notmuch/0.29.1+93~g67ed7df (https://notmuchmail.org) Emacs/26.3 (x86_64-pc-linux-gnu) Date: Thu, 15 Jul 2021 12:33:04 +0200 Message-ID: <8735sfamu7.fsf@dem-tschwing-1.ger.mentorg.com> MIME-Version: 1.0 Content-Type: text/plain; charset="utf-8" Content-Transfer-Encoding: quoted-printable X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: SVR-IES-MBX-07.mgc.mentorg.com (139.181.222.7) To svr-ies-mbx-01.mgc.mentorg.com (139.181.222.1) X-Spam-Status: No, score=-12.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.4 X-Spam-Checker-Version: SpamAssassin 3.4.4 (2020-01-24) 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: , X-List-Received-Date: Thu, 15 Jul 2021 10:33:18 -0000 Hi! On 2021-07-02T09:15:27+0200, Richard Biener via Gcc-patches wrote: > On Thu, Jul 1, 2021 at 5:17 PM Hafiz Abid Qadeer = wrote: >> >> Currently, if we look at the debug information for offload kernel >> regions, it looks something like this: >> >> void foo (void) >> { >> #pragma acc kernels >> { >> >> } >> } >> >> DW_TAG_compile_unit >> DW_AT_name ("") >> >> DW_TAG_subprogram // notional parent function (foo) with no code range >> >> DW_TAG_subprogram // offload function foo._omp_fn.0 >> >> There is an artificial compile unit. It contains a parent subprogram whi= ch >> has the offload function as its child. The parent function makes sense = in >> host code where it actually exists and does have an address range. But i= n >> offload code, it does not exist and neither the generated dwarf has an >> address range for this function. >> >> When debugger read the dwarf for offload code, they see a function with = no >> address range and discard it alongwith its children which include offloa= d >> function. This results in a poor debug experience of offload code. >> >> This patch tries to solve this problem by making offload kernels childre= n of >> "artifical" compile unit instead of a non existent parent function. This >> not only improves debug experience but also reflects the reality better >> in debug info. >> >> Patch was tested on x86_64 with amdgcn offload. Debug behavior was >> tested with rocgdb. > > The proper fix is to reflect this in the functions declaration which curr= ently > will have a DECL_CONTEXT of the containing function. That could be > done either on the host as well or alternatively at the time we offload > the "child" but not the parent. Does that mean adding a (very simple) new pass in the offloading compilation pipeline, conditionalizing this 'DECL_CONTEXT' modification under '#ifdef ACCEL_COMPILER'? See 'gcc/omp-offload.c:pass_omp_target_link' for a simple example. Should that be placed at the beginning of the offloading pipeline, thus before 'pass_oacc_device_lower' (see 'gcc/passes.def'), or doesn't matter where, I suppose? Please cross-reference 'gcc/omp-low.c:create_omp_child_function', 'gcc/omp-expand.c:adjust_context_and_scope', and the new pass, assuming these are the relevant pieces here? > Note that the "parent" should be abstract but I don't think dwarf has a > way to express a fully abstract parent of a concrete instance child - or > at least how GCC expresses this causes consumers to "misinterpret" > that. I wonder if adding a DW_AT_declaration to the late DWARF > emitted "parent" would fix things as well here? (I suppose not, Abid?) Gr=C3=BC=C3=9Fe Thomas >> gcc/ >> >> * gcc/dwarf2out.c (notional_parents_list): New file variable. >> (gen_subprogram_die): Record offload kernel functions in >> notional_parents_list. >> (fixup_notional_parents): New function. >> (dwarf2out_finish): Call fixup_notional_parents. >> (dwarf2out_c_finalize): Reset notional_parents_list. >> --- >> gcc/dwarf2out.c | 68 +++++++++++++++++++++++++++++++++++++++++++++++-- >> 1 file changed, 66 insertions(+), 2 deletions(-) >> >> diff --git a/gcc/dwarf2out.c b/gcc/dwarf2out.c >> index 80acf165fee..769bb7fc4a8 100644 >> --- a/gcc/dwarf2out.c >> +++ b/gcc/dwarf2out.c >> @@ -3506,6 +3506,11 @@ static GTY(()) limbo_die_node *limbo_die_list; >> DW_AT_{,MIPS_}linkage_name once their DECL_ASSEMBLER_NAMEs are set. = */ >> static GTY(()) limbo_die_node *deferred_asm_name; >> >> +/* A list of DIEs which represent parents of nested offload kernels. T= hese >> + functions exist on the host side but not in the offloed code. But t= hey >> + still show up as parent of the ofload kernels in DWARF. */ >> +static GTY(()) limbo_die_node *notional_parents_list; >> + >> struct dwarf_file_hasher : ggc_ptr_hash >> { >> typedef const char *compare_type; >> @@ -23652,8 +23657,23 @@ gen_subprogram_die (tree decl, dw_die_ref conte= xt_die) >> if (fde->dw_fde_begin) >> { >> /* We have already generated the labels. */ >> - add_AT_low_high_pc (subr_die, fde->dw_fde_begin, >> - fde->dw_fde_end, false); >> + add_AT_low_high_pc (subr_die, fde->dw_fde_begin, >> + fde->dw_fde_end, false); >> + >> + /* Offload kernel functions are nested within a parent func= tion >> + that doesn't actually exist in the offload object. GDB >> + will ignore the function and everything nested within it= as >> + the function does not have an address range. We mark th= e >> + parent functions here and will later fix them. */ >> + if (lookup_attribute ("omp target entrypoint", >> + DECL_ATTRIBUTES (decl))) >> + { >> + limbo_die_node *node =3D ggc_cleared_alloc (); >> + node->die =3D subr_die->die_parent; >> + node->created_for =3D decl; >> + node->next =3D notional_parents_list; >> + notional_parents_list =3D node; >> + } >> } >> else >> { >> @@ -31881,6 +31901,46 @@ flush_limbo_die_list (void) >> } >> } >> >> +/* Fixup notional parent function (which does not actually exist) so th= at >> + a function with no address range is not parent of a function *with* = address >> + ranges. Otherwise debugger see the parent function without code ran= ge >> + and discards it along with its children which here include function >> + which have address range. >> + >> + Typically this occurs when we have an offload kernel, where the pare= nt >> + function only exists in the host-side portion of the code. */ >> + >> +static void >> +fixup_notional_parents (void) >> +{ >> + limbo_die_node *node; >> + >> + for (node =3D notional_parents_list; node; node =3D node->next) >> + { >> + dw_die_ref notional_parent =3D node->die; >> + /* The dwarf at this moment looks like this >> + DW_TAG_compile_unit >> + DW_AT_name ("") >> + >> + DW_TAG_subprogram // parent function with no code range >> + >> + DW_TAG_subprogram // offload function 1 >> + ... >> + DW_TAG_subprogram // offload function n >> + Our aim is to make offload function children of CU. */ >> + if (notional_parent >> + && notional_parent->die_tag =3D=3D DW_TAG_subprogram >> + && !(get_AT (notional_parent, DW_AT_low_pc) >> + || get_AT (notional_parent, DW_AT_ranges))) >> + >> + { >> + dw_die_ref cu =3D notional_parent->die_parent; >> + if (cu && cu->die_tag =3D=3D DW_TAG_compile_unit) >> + reparent_child (notional_parent->die_child, cu); >> + } >> + } >> +} >> + >> /* Reset DIEs so we can output them again. */ >> >> static void >> @@ -31938,6 +31998,9 @@ dwarf2out_finish (const char *filename) >> /* Flush out any latecomers to the limbo party. */ >> flush_limbo_die_list (); >> >> + /* Sort out notional parents of offloaded kernel. */ >> + fixup_notional_parents (); >> + >> if (inline_entry_data_table) >> gcc_assert (inline_entry_data_table->is_empty ()); >> >> @@ -32994,6 +33057,7 @@ dwarf2out_c_finalize (void) >> single_comp_unit_die =3D NULL; >> comdat_type_list =3D NULL; >> limbo_die_list =3D NULL; >> + notional_parents_list =3D NULL; >> file_table =3D NULL; >> decl_die_table =3D NULL; >> common_block_die_table =3D NULL; >> -- >> 2.25.1 >> ----------------- 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