| Message ID | 20210701151657.935006-1-abidh@codesourcery.com |
|---|---|
| State | New |
| Headers | show |
| Series | [DWARF] Fix hierarchy of debug information for offload kernels. | expand |
On Thu, Jul 1, 2021 at 5:17 PM Hafiz Abid Qadeer <abidh@codesourcery.com> 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 ("<artificial>") > > 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 which > 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 in > 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 offload > function. This results in a poor debug experience of offload code. > > This patch tries to solve this problem by making offload kernels children 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 currently 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. 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? Richard. > 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. These > + functions exist on the host side but not in the offloed code. But they > + 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<dwarf_file_data> > { > typedef const char *compare_type; > @@ -23652,8 +23657,23 @@ gen_subprogram_die (tree decl, dw_die_ref context_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 function > + 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 the > + parent functions here and will later fix them. */ > + if (lookup_attribute ("omp target entrypoint", > + DECL_ATTRIBUTES (decl))) > + { > + limbo_die_node *node = ggc_cleared_alloc<limbo_die_node> (); > + node->die = subr_die->die_parent; > + node->created_for = decl; > + node->next = notional_parents_list; > + notional_parents_list = node; > + } > } > else > { > @@ -31881,6 +31901,46 @@ flush_limbo_die_list (void) > } > } > > +/* Fixup notional parent function (which does not actually exist) so that > + a function with no address range is not parent of a function *with* address > + ranges. Otherwise debugger see the parent function without code range > + 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 parent > + function only exists in the host-side portion of the code. */ > + > +static void > +fixup_notional_parents (void) > +{ > + limbo_die_node *node; > + > + for (node = notional_parents_list; node; node = node->next) > + { > + dw_die_ref notional_parent = node->die; > + /* The dwarf at this moment looks like this > + DW_TAG_compile_unit > + DW_AT_name ("<artificial>") > + > + 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 == DW_TAG_subprogram > + && !(get_AT (notional_parent, DW_AT_low_pc) > + || get_AT (notional_parent, DW_AT_ranges))) > + > + { > + dw_die_ref cu = notional_parent->die_parent; > + if (cu && cu->die_tag == 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 = NULL; > comdat_type_list = NULL; > limbo_die_list = NULL; > + notional_parents_list = NULL; > file_table = NULL; > decl_die_table = NULL; > common_block_die_table = NULL; > -- > 2.25.1 >
Hi! On 2021-07-02T09:15:27+0200, Richard Biener via Gcc-patches <gcc-patches@gcc.gnu.org> wrote: > On Thu, Jul 1, 2021 at 5:17 PM Hafiz Abid Qadeer <abidh@codesourcery.com> 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 ("<artificial>") >> >> 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 which >> 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 in >> 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 offload >> function. This results in a poor debug experience of offload code. >> >> This patch tries to solve this problem by making offload kernels children 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 currently > 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üße 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. These >> + functions exist on the host side but not in the offloed code. But they >> + 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<dwarf_file_data> >> { >> typedef const char *compare_type; >> @@ -23652,8 +23657,23 @@ gen_subprogram_die (tree decl, dw_die_ref context_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 function >> + 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 the >> + parent functions here and will later fix them. */ >> + if (lookup_attribute ("omp target entrypoint", >> + DECL_ATTRIBUTES (decl))) >> + { >> + limbo_die_node *node = ggc_cleared_alloc<limbo_die_node> (); >> + node->die = subr_die->die_parent; >> + node->created_for = decl; >> + node->next = notional_parents_list; >> + notional_parents_list = node; >> + } >> } >> else >> { >> @@ -31881,6 +31901,46 @@ flush_limbo_die_list (void) >> } >> } >> >> +/* Fixup notional parent function (which does not actually exist) so that >> + a function with no address range is not parent of a function *with* address >> + ranges. Otherwise debugger see the parent function without code range >> + 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 parent >> + function only exists in the host-side portion of the code. */ >> + >> +static void >> +fixup_notional_parents (void) >> +{ >> + limbo_die_node *node; >> + >> + for (node = notional_parents_list; node; node = node->next) >> + { >> + dw_die_ref notional_parent = node->die; >> + /* The dwarf at this moment looks like this >> + DW_TAG_compile_unit >> + DW_AT_name ("<artificial>") >> + >> + 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 == DW_TAG_subprogram >> + && !(get_AT (notional_parent, DW_AT_low_pc) >> + || get_AT (notional_parent, DW_AT_ranges))) >> + >> + { >> + dw_die_ref cu = notional_parent->die_parent; >> + if (cu && cu->die_tag == 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 = NULL; >> comdat_type_list = NULL; >> limbo_die_list = NULL; >> + notional_parents_list = NULL; >> file_table = NULL; >> decl_die_table = NULL; >> common_block_die_table = NULL; >> -- >> 2.25.1 >> ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
On 15/07/2021 11:33, Thomas Schwinge wrote: > >> 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?) > Yes, adding DW_AT_declaration does not fix the problem.
On Thu, Jul 15, 2021 at 12:35 PM Hafiz Abid Qadeer <abid_qadeer@mentor.com> wrote: > > On 15/07/2021 11:33, Thomas Schwinge wrote: > > > >> 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?) > > > > Yes, adding DW_AT_declaration does not fix the problem. Does emitting DW_TAG_compile_unit DW_AT_name ("<artificial>") DW_TAG_subprogram // notional parent function (foo) with no code range DW_AT_declaration 1 a: DW_TAG_subprogram // offload function foo._omp_fn.0 DW_AT_declaration 1 DW_TAG_subprogram // offload function DW_AT_abstract_origin a ... do the trick? The following would do this, flattening function definitions for the concrete copies: diff --git a/gcc/dwarf2out.c b/gcc/dwarf2out.c index 82783c4968b..a9c8bc43e88 100644 --- a/gcc/dwarf2out.c +++ b/gcc/dwarf2out.c @@ -6076,6 +6076,11 @@ maybe_create_die_with_external_ref (tree decl) /* Peel types in the context stack. */ while (ctx && TYPE_P (ctx)) ctx = TYPE_CONTEXT (ctx); + /* For functions peel the context up to namespace/TU scope. The abstract + copies reveal the true nesting. */ + if (TREE_CODE (decl) == FUNCTION_DECL) + while (ctx && TREE_CODE (ctx) == FUNCTION_DECL) + ctx = DECL_CONTEXT (ctx); /* Likewise namespaces in case we do not want to emit DIEs for them. */ if (debug_info_level <= DINFO_LEVEL_TERSE) while (ctx && TREE_CODE (ctx) == NAMESPACE_DECL) @@ -6099,8 +6104,7 @@ maybe_create_die_with_external_ref (tree decl) /* Leave function local entities parent determination to when we process scope vars. */ ; - else - parent = lookup_decl_die (ctx); + parent = lookup_decl_die (ctx); } else /* In some cases the FEs fail to set DECL_CONTEXT properly. > > -- > Hafiz Abid Qadeer > Mentor, a Siemens Business
On 15/07/2021 13:09, Richard Biener wrote: > On Thu, Jul 15, 2021 at 12:35 PM Hafiz Abid Qadeer > <abid_qadeer@mentor.com> wrote: >> >> On 15/07/2021 11:33, Thomas Schwinge wrote: >>> >>>> 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?) >>> >> >> Yes, adding DW_AT_declaration does not fix the problem. > > Does emitting > > DW_TAG_compile_unit > DW_AT_name ("<artificial>") > > DW_TAG_subprogram // notional parent function (foo) with no code range > DW_AT_declaration 1 > a: DW_TAG_subprogram // offload function foo._omp_fn.0 > DW_AT_declaration 1 > > DW_TAG_subprogram // offload function > DW_AT_abstract_origin a > ... > > do the trick? The following would do this, flattening function definitions > for the concrete copies: > > diff --git a/gcc/dwarf2out.c b/gcc/dwarf2out.c > index 82783c4968b..a9c8bc43e88 100644 > --- a/gcc/dwarf2out.c > +++ b/gcc/dwarf2out.c > @@ -6076,6 +6076,11 @@ maybe_create_die_with_external_ref (tree decl) > /* Peel types in the context stack. */ > while (ctx && TYPE_P (ctx)) > ctx = TYPE_CONTEXT (ctx); > + /* For functions peel the context up to namespace/TU scope. The abstract > + copies reveal the true nesting. */ > + if (TREE_CODE (decl) == FUNCTION_DECL) > + while (ctx && TREE_CODE (ctx) == FUNCTION_DECL) > + ctx = DECL_CONTEXT (ctx); > /* Likewise namespaces in case we do not want to emit DIEs for them. */ > if (debug_info_level <= DINFO_LEVEL_TERSE) > while (ctx && TREE_CODE (ctx) == NAMESPACE_DECL) > @@ -6099,8 +6104,7 @@ maybe_create_die_with_external_ref (tree decl) > /* Leave function local entities parent determination to when > we process scope vars. */ > ; > - else > - parent = lookup_decl_die (ctx); > + parent = lookup_decl_die (ctx); > } > else > /* In some cases the FEs fail to set DECL_CONTEXT properly. > Thanks. This solves the problem. Only the first hunk was required. Second hunk actually causes an ICE when TREE_CODE (ctx) == BLOCK. OK to commit the attached patch?
On Fri, Jul 16, 2021 at 10:23 PM Hafiz Abid Qadeer <abid_qadeer@mentor.com> wrote: > > On 15/07/2021 13:09, Richard Biener wrote: > > On Thu, Jul 15, 2021 at 12:35 PM Hafiz Abid Qadeer > > <abid_qadeer@mentor.com> wrote: > >> > >> On 15/07/2021 11:33, Thomas Schwinge wrote: > >>> > >>>> 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?) > >>> > >> > >> Yes, adding DW_AT_declaration does not fix the problem. > > > > Does emitting > > > > DW_TAG_compile_unit > > DW_AT_name ("<artificial>") > > > > DW_TAG_subprogram // notional parent function (foo) with no code range > > DW_AT_declaration 1 > > a: DW_TAG_subprogram // offload function foo._omp_fn.0 > > DW_AT_declaration 1 > > > > DW_TAG_subprogram // offload function > > DW_AT_abstract_origin a > > ... > > > > do the trick? The following would do this, flattening function definitions > > for the concrete copies: > > > > diff --git a/gcc/dwarf2out.c b/gcc/dwarf2out.c > > index 82783c4968b..a9c8bc43e88 100644 > > --- a/gcc/dwarf2out.c > > +++ b/gcc/dwarf2out.c > > @@ -6076,6 +6076,11 @@ maybe_create_die_with_external_ref (tree decl) > > /* Peel types in the context stack. */ > > while (ctx && TYPE_P (ctx)) > > ctx = TYPE_CONTEXT (ctx); > > + /* For functions peel the context up to namespace/TU scope. The abstract > > + copies reveal the true nesting. */ > > + if (TREE_CODE (decl) == FUNCTION_DECL) > > + while (ctx && TREE_CODE (ctx) == FUNCTION_DECL) > > + ctx = DECL_CONTEXT (ctx); > > /* Likewise namespaces in case we do not want to emit DIEs for them. */ > > if (debug_info_level <= DINFO_LEVEL_TERSE) > > while (ctx && TREE_CODE (ctx) == NAMESPACE_DECL) > > @@ -6099,8 +6104,7 @@ maybe_create_die_with_external_ref (tree decl) > > /* Leave function local entities parent determination to when > > we process scope vars. */ > > ; > > - else > > - parent = lookup_decl_die (ctx); > > + parent = lookup_decl_die (ctx); > > } > > else > > /* In some cases the FEs fail to set DECL_CONTEXT properly. > > > > Thanks. This solves the problem. Only the first hunk was required. Second hunk > actually causes an ICE when TREE_CODE (ctx) == BLOCK. > OK to commit the attached patch? I think we need to merge the TYPE_P peeling and FUNCTION_DECL peeling into one loop since I suppose we can have a nested function in class scope. So sth like diff --git a/gcc/dwarf2out.c b/gcc/dwarf2out.c index 82783c4968b..61228410b51 100644 --- a/gcc/dwarf2out.c +++ b/gcc/dwarf2out.c @@ -6073,8 +6073,12 @@ maybe_create_die_with_external_ref (tree decl) } else ctx = DECL_CONTEXT (decl); - /* Peel types in the context stack. */ - while (ctx && TYPE_P (ctx)) + /* Peel types in the context stack. For functions peel the context up + to namespace/TU scope. The abstract copies reveal the true nesting. */ + while (ctx + && (TYPE_P (ctx) + || (TREE_CODE (decl) == FUNCTION_DECL + && TREE_CODE (ctx) == FUNCTION_DECL))) ctx = TYPE_CONTEXT (ctx); /* Likewise namespaces in case we do not want to emit DIEs for them. */ if (debug_info_level <= DINFO_LEVEL_TERSE) if that works it's OK. Can you run it on the gdb testsuite with -flto added as well please (you need to do before/after comparison since IIRC adding -flto will add a few fails). Thanks, Richard.
On 19/07/2021 11:45, Richard Biener wrote: > On Fri, Jul 16, 2021 at 10:23 PM Hafiz Abid Qadeer > <abid_qadeer@mentor.com> wrote: >> >> On 15/07/2021 13:09, Richard Biener wrote: >>> On Thu, Jul 15, 2021 at 12:35 PM Hafiz Abid Qadeer >>> <abid_qadeer@mentor.com> wrote: >>>> >>>> On 15/07/2021 11:33, Thomas Schwinge wrote: >>>>> >>>>>> 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?) >>>>> >>>> >>>> Yes, adding DW_AT_declaration does not fix the problem. >>> >>> Does emitting >>> >>> DW_TAG_compile_unit >>> DW_AT_name ("<artificial>") >>> >>> DW_TAG_subprogram // notional parent function (foo) with no code range >>> DW_AT_declaration 1 >>> a: DW_TAG_subprogram // offload function foo._omp_fn.0 >>> DW_AT_declaration 1 >>> >>> DW_TAG_subprogram // offload function >>> DW_AT_abstract_origin a >>> ... >>> >>> do the trick? The following would do this, flattening function definitions >>> for the concrete copies: >>> >>> diff --git a/gcc/dwarf2out.c b/gcc/dwarf2out.c >>> index 82783c4968b..a9c8bc43e88 100644 >>> --- a/gcc/dwarf2out.c >>> +++ b/gcc/dwarf2out.c >>> @@ -6076,6 +6076,11 @@ maybe_create_die_with_external_ref (tree decl) >>> /* Peel types in the context stack. */ >>> while (ctx && TYPE_P (ctx)) >>> ctx = TYPE_CONTEXT (ctx); >>> + /* For functions peel the context up to namespace/TU scope. The abstract >>> + copies reveal the true nesting. */ >>> + if (TREE_CODE (decl) == FUNCTION_DECL) >>> + while (ctx && TREE_CODE (ctx) == FUNCTION_DECL) >>> + ctx = DECL_CONTEXT (ctx); >>> /* Likewise namespaces in case we do not want to emit DIEs for them. */ >>> if (debug_info_level <= DINFO_LEVEL_TERSE) >>> while (ctx && TREE_CODE (ctx) == NAMESPACE_DECL) >>> @@ -6099,8 +6104,7 @@ maybe_create_die_with_external_ref (tree decl) >>> /* Leave function local entities parent determination to when >>> we process scope vars. */ >>> ; >>> - else >>> - parent = lookup_decl_die (ctx); >>> + parent = lookup_decl_die (ctx); >>> } >>> else >>> /* In some cases the FEs fail to set DECL_CONTEXT properly. >>> >> >> Thanks. This solves the problem. Only the first hunk was required. Second hunk >> actually causes an ICE when TREE_CODE (ctx) == BLOCK. >> OK to commit the attached patch? > > I think we need to merge the TYPE_P peeling and FUNCTION_DECL peeling into > one loop since I suppose we can have a nested function in class scope. > So sth like > > diff --git a/gcc/dwarf2out.c b/gcc/dwarf2out.c > index 82783c4968b..61228410b51 100644 > --- a/gcc/dwarf2out.c > +++ b/gcc/dwarf2out.c > @@ -6073,8 +6073,12 @@ maybe_create_die_with_external_ref (tree decl) > } > else > ctx = DECL_CONTEXT (decl); > - /* Peel types in the context stack. */ > - while (ctx && TYPE_P (ctx)) > + /* Peel types in the context stack. For functions peel the context up > + to namespace/TU scope. The abstract copies reveal the true nesting. */ > + while (ctx > + && (TYPE_P (ctx) > + || (TREE_CODE (decl) == FUNCTION_DECL > + && TREE_CODE (ctx) == FUNCTION_DECL))) > ctx = TYPE_CONTEXT (ctx); > /* Likewise namespaces in case we do not want to emit DIEs for them. */ > if (debug_info_level <= DINFO_LEVEL_TERSE) > This causes an ICE, internal compiler error: tree check: expected class 'type', have 'declaration' (function_decl) Did you intend something like this: diff --git a/gcc/dwarf2out.c b/gcc/dwarf2out.c index 561f8b23517..c61f0041fba 100644 --- a/gcc/dwarf2out.c +++ b/gcc/dwarf2out.c @@ -6121,3 +6121,8 @@ maybe_create_die_with_external_ref (tree decl) - /* Peel types in the context stack. */ - while (ctx && TYPE_P (ctx)) - ctx = TYPE_CONTEXT (ctx); + /* Peel types in the context stack. For functions peel the context up + to namespace/TU scope. The abstract copies reveal the true nesting. */ + while (ctx + && (TYPE_P (ctx) + || (TREE_CODE (decl) == FUNCTION_DECL + && TREE_CODE (ctx) == FUNCTION_DECL))) + ctx = TYPE_P (ctx) ? TYPE_CONTEXT (ctx) : DECL_CONTEXT (ctx); + > if that works it's OK. Can you run it on the gdb testsuite with -flto added > as well please (you need to do before/after comparison since IIRC adding > -flto will add a few fails). > > Thanks, > Richard. >
On July 19, 2021 6:13:40 PM GMT+02:00, Hafiz Abid Qadeer <abid_qadeer@mentor.com> wrote: >On 19/07/2021 11:45, Richard Biener wrote: >> On Fri, Jul 16, 2021 at 10:23 PM Hafiz Abid Qadeer >> <abid_qadeer@mentor.com> wrote: >>> >>> On 15/07/2021 13:09, Richard Biener wrote: >>>> On Thu, Jul 15, 2021 at 12:35 PM Hafiz Abid Qadeer >>>> <abid_qadeer@mentor.com> wrote: >>>>> >>>>> On 15/07/2021 11:33, Thomas Schwinge wrote: >>>>>> >>>>>>> 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?) >>>>>> >>>>> >>>>> Yes, adding DW_AT_declaration does not fix the problem. >>>> >>>> Does emitting >>>> >>>> DW_TAG_compile_unit >>>> DW_AT_name ("<artificial>") >>>> >>>> DW_TAG_subprogram // notional parent function (foo) with no code >range >>>> DW_AT_declaration 1 >>>> a: DW_TAG_subprogram // offload function foo._omp_fn.0 >>>> DW_AT_declaration 1 >>>> >>>> DW_TAG_subprogram // offload function >>>> DW_AT_abstract_origin a >>>> ... >>>> >>>> do the trick? The following would do this, flattening function >definitions >>>> for the concrete copies: >>>> >>>> diff --git a/gcc/dwarf2out.c b/gcc/dwarf2out.c >>>> index 82783c4968b..a9c8bc43e88 100644 >>>> --- a/gcc/dwarf2out.c >>>> +++ b/gcc/dwarf2out.c >>>> @@ -6076,6 +6076,11 @@ maybe_create_die_with_external_ref (tree >decl) >>>> /* Peel types in the context stack. */ >>>> while (ctx && TYPE_P (ctx)) >>>> ctx = TYPE_CONTEXT (ctx); >>>> + /* For functions peel the context up to namespace/TU scope. The >abstract >>>> + copies reveal the true nesting. */ >>>> + if (TREE_CODE (decl) == FUNCTION_DECL) >>>> + while (ctx && TREE_CODE (ctx) == FUNCTION_DECL) >>>> + ctx = DECL_CONTEXT (ctx); >>>> /* Likewise namespaces in case we do not want to emit DIEs for >them. */ >>>> if (debug_info_level <= DINFO_LEVEL_TERSE) >>>> while (ctx && TREE_CODE (ctx) == NAMESPACE_DECL) >>>> @@ -6099,8 +6104,7 @@ maybe_create_die_with_external_ref (tree >decl) >>>> /* Leave function local entities parent determination to >when >>>> we process scope vars. */ >>>> ; >>>> - else >>>> - parent = lookup_decl_die (ctx); >>>> + parent = lookup_decl_die (ctx); >>>> } >>>> else >>>> /* In some cases the FEs fail to set DECL_CONTEXT properly. >>>> >>> >>> Thanks. This solves the problem. Only the first hunk was required. >Second hunk >>> actually causes an ICE when TREE_CODE (ctx) == BLOCK. >>> OK to commit the attached patch? >> >> I think we need to merge the TYPE_P peeling and FUNCTION_DECL peeling >into >> one loop since I suppose we can have a nested function in class >scope. >> So sth like >> >> diff --git a/gcc/dwarf2out.c b/gcc/dwarf2out.c >> index 82783c4968b..61228410b51 100644 >> --- a/gcc/dwarf2out.c >> +++ b/gcc/dwarf2out.c >> @@ -6073,8 +6073,12 @@ maybe_create_die_with_external_ref (tree decl) >> } >> else >> ctx = DECL_CONTEXT (decl); >> - /* Peel types in the context stack. */ >> - while (ctx && TYPE_P (ctx)) >> + /* Peel types in the context stack. For functions peel the >context up >> + to namespace/TU scope. The abstract copies reveal the true >nesting. */ >> + while (ctx >> + && (TYPE_P (ctx) >> + || (TREE_CODE (decl) == FUNCTION_DECL >> + && TREE_CODE (ctx) == FUNCTION_DECL))) >> ctx = TYPE_CONTEXT (ctx); >> /* Likewise namespaces in case we do not want to emit DIEs for >them. */ >> if (debug_info_level <= DINFO_LEVEL_TERSE) >> >This causes an ICE, >internal compiler error: tree check: expected class 'type', have >'declaration' (function_decl) > >Did you intend something like this: > >diff --git a/gcc/dwarf2out.c b/gcc/dwarf2out.c >index 561f8b23517..c61f0041fba 100644 >--- a/gcc/dwarf2out.c >+++ b/gcc/dwarf2out.c >@@ -6121,3 +6121,8 @@ maybe_create_die_with_external_ref (tree decl) >- /* Peel types in the context stack. */ >- while (ctx && TYPE_P (ctx)) >- ctx = TYPE_CONTEXT (ctx); >+ /* Peel types in the context stack. For functions peel the context >up >+ to namespace/TU scope. The abstract copies reveal the true >nesting. */ >+ while (ctx >+ && (TYPE_P (ctx) >+ || (TREE_CODE (decl) == FUNCTION_DECL >+ && TREE_CODE (ctx) == FUNCTION_DECL))) >+ ctx = TYPE_P (ctx) ? TYPE_CONTEXT (ctx) : DECL_CONTEXT (ctx); >+ Yes, of course. > >> if that works it's OK. Can you run it on the gdb testsuite with >-flto added >> as well please (you need to do before/after comparison since IIRC >adding >> -flto will add a few fails). >> >> Thanks, >> Richard. >>
On 19/07/2021 17:41, Richard Biener wrote: > On July 19, 2021 6:13:40 PM GMT+02:00, Hafiz Abid Qadeer <abid_qadeer@mentor.com> wrote: >> On 19/07/2021 11:45, Richard Biener wrote: >>> On Fri, Jul 16, 2021 at 10:23 PM Hafiz Abid Qadeer >>> <abid_qadeer@mentor.com> wrote: >>>> >>>> On 15/07/2021 13:09, Richard Biener wrote: >>>>> On Thu, Jul 15, 2021 at 12:35 PM Hafiz Abid Qadeer >>>>> <abid_qadeer@mentor.com> wrote: >>>>>> >>>>>> On 15/07/2021 11:33, Thomas Schwinge wrote: >>>>>>> >>>>>>>> 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?) >>>>>>> >>>>>> >>>>>> Yes, adding DW_AT_declaration does not fix the problem. >>>>> >>>>> Does emitting >>>>> >>>>> DW_TAG_compile_unit >>>>> DW_AT_name ("<artificial>") >>>>> >>>>> DW_TAG_subprogram // notional parent function (foo) with no code >> range >>>>> DW_AT_declaration 1 >>>>> a: DW_TAG_subprogram // offload function foo._omp_fn.0 >>>>> DW_AT_declaration 1 >>>>> >>>>> DW_TAG_subprogram // offload function >>>>> DW_AT_abstract_origin a >>>>> ... >>>>> >>>>> do the trick? The following would do this, flattening function >> definitions >>>>> for the concrete copies: >>>>> >>>>> diff --git a/gcc/dwarf2out.c b/gcc/dwarf2out.c >>>>> index 82783c4968b..a9c8bc43e88 100644 >>>>> --- a/gcc/dwarf2out.c >>>>> +++ b/gcc/dwarf2out.c >>>>> @@ -6076,6 +6076,11 @@ maybe_create_die_with_external_ref (tree >> decl) >>>>> /* Peel types in the context stack. */ >>>>> while (ctx && TYPE_P (ctx)) >>>>> ctx = TYPE_CONTEXT (ctx); >>>>> + /* For functions peel the context up to namespace/TU scope. The >> abstract >>>>> + copies reveal the true nesting. */ >>>>> + if (TREE_CODE (decl) == FUNCTION_DECL) >>>>> + while (ctx && TREE_CODE (ctx) == FUNCTION_DECL) >>>>> + ctx = DECL_CONTEXT (ctx); >>>>> /* Likewise namespaces in case we do not want to emit DIEs for >> them. */ >>>>> if (debug_info_level <= DINFO_LEVEL_TERSE) >>>>> while (ctx && TREE_CODE (ctx) == NAMESPACE_DECL) >>>>> @@ -6099,8 +6104,7 @@ maybe_create_die_with_external_ref (tree >> decl) >>>>> /* Leave function local entities parent determination to >> when >>>>> we process scope vars. */ >>>>> ; >>>>> - else >>>>> - parent = lookup_decl_die (ctx); >>>>> + parent = lookup_decl_die (ctx); >>>>> } >>>>> else >>>>> /* In some cases the FEs fail to set DECL_CONTEXT properly. >>>>> >>>> >>>> Thanks. This solves the problem. Only the first hunk was required. >> Second hunk >>>> actually causes an ICE when TREE_CODE (ctx) == BLOCK. >>>> OK to commit the attached patch? >>> >>> I think we need to merge the TYPE_P peeling and FUNCTION_DECL peeling >> into >>> one loop since I suppose we can have a nested function in class >> scope. >>> So sth like >>> >>> diff --git a/gcc/dwarf2out.c b/gcc/dwarf2out.c >>> index 82783c4968b..61228410b51 100644 >>> --- a/gcc/dwarf2out.c >>> +++ b/gcc/dwarf2out.c >>> @@ -6073,8 +6073,12 @@ maybe_create_die_with_external_ref (tree decl) >>> } >>> else >>> ctx = DECL_CONTEXT (decl); >>> - /* Peel types in the context stack. */ >>> - while (ctx && TYPE_P (ctx)) >>> + /* Peel types in the context stack. For functions peel the >> context up >>> + to namespace/TU scope. The abstract copies reveal the true >> nesting. */ >>> + while (ctx >>> + && (TYPE_P (ctx) >>> + || (TREE_CODE (decl) == FUNCTION_DECL >>> + && TREE_CODE (ctx) == FUNCTION_DECL))) >>> ctx = TYPE_CONTEXT (ctx); >>> /* Likewise namespaces in case we do not want to emit DIEs for >> them. */ >>> if (debug_info_level <= DINFO_LEVEL_TERSE) >>> >> This causes an ICE, >> internal compiler error: tree check: expected class 'type', have >> 'declaration' (function_decl) >> >> Did you intend something like this: >> >> diff --git a/gcc/dwarf2out.c b/gcc/dwarf2out.c >> index 561f8b23517..c61f0041fba 100644 >> --- a/gcc/dwarf2out.c >> +++ b/gcc/dwarf2out.c >> @@ -6121,3 +6121,8 @@ maybe_create_die_with_external_ref (tree decl) >> - /* Peel types in the context stack. */ >> - while (ctx && TYPE_P (ctx)) >> - ctx = TYPE_CONTEXT (ctx); >> + /* Peel types in the context stack. For functions peel the context >> up >> + to namespace/TU scope. The abstract copies reveal the true >> nesting. */ >> + while (ctx >> + && (TYPE_P (ctx) >> + || (TREE_CODE (decl) == FUNCTION_DECL >> + && TREE_CODE (ctx) == FUNCTION_DECL))) >> + ctx = TYPE_P (ctx) ? TYPE_CONTEXT (ctx) : DECL_CONTEXT (ctx); >> + > > Yes, of course. > >> >>> if that works it's OK. Can you run it on the gdb testsuite with >> -flto added >>> as well please (you need to do before/after comparison since IIRC >> adding >>> -flto will add a few fails). GDB testsuite shows some extra fails which mainly happen because testcase assumes that you can access the local variable of enclosing function from the nested function (or omp parallel region). After this change, the nested functions are no longer children of the enclosing function so those tests fail. The problem that prompted this patch happened for parent function that did not have a code range i.e a notional parent. I was wondering if we should update the ctx only for such parents instead of all function as we did above. Thanks,
On Wed, Jul 21, 2021 at 7:55 PM Hafiz Abid Qadeer <abid_qadeer@mentor.com> wrote: > > On 19/07/2021 17:41, Richard Biener wrote: > > On July 19, 2021 6:13:40 PM GMT+02:00, Hafiz Abid Qadeer <abid_qadeer@mentor.com> wrote: > >> On 19/07/2021 11:45, Richard Biener wrote: > >>> On Fri, Jul 16, 2021 at 10:23 PM Hafiz Abid Qadeer > >>> <abid_qadeer@mentor.com> wrote: > >>>> > >>>> On 15/07/2021 13:09, Richard Biener wrote: > >>>>> On Thu, Jul 15, 2021 at 12:35 PM Hafiz Abid Qadeer > >>>>> <abid_qadeer@mentor.com> wrote: > >>>>>> > >>>>>> On 15/07/2021 11:33, Thomas Schwinge wrote: > >>>>>>> > >>>>>>>> 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?) > >>>>>>> > >>>>>> > >>>>>> Yes, adding DW_AT_declaration does not fix the problem. > >>>>> > >>>>> Does emitting > >>>>> > >>>>> DW_TAG_compile_unit > >>>>> DW_AT_name ("<artificial>") > >>>>> > >>>>> DW_TAG_subprogram // notional parent function (foo) with no code > >> range > >>>>> DW_AT_declaration 1 > >>>>> a: DW_TAG_subprogram // offload function foo._omp_fn.0 > >>>>> DW_AT_declaration 1 > >>>>> > >>>>> DW_TAG_subprogram // offload function > >>>>> DW_AT_abstract_origin a > >>>>> ... > >>>>> > >>>>> do the trick? The following would do this, flattening function > >> definitions > >>>>> for the concrete copies: > >>>>> > >>>>> diff --git a/gcc/dwarf2out.c b/gcc/dwarf2out.c > >>>>> index 82783c4968b..a9c8bc43e88 100644 > >>>>> --- a/gcc/dwarf2out.c > >>>>> +++ b/gcc/dwarf2out.c > >>>>> @@ -6076,6 +6076,11 @@ maybe_create_die_with_external_ref (tree > >> decl) > >>>>> /* Peel types in the context stack. */ > >>>>> while (ctx && TYPE_P (ctx)) > >>>>> ctx = TYPE_CONTEXT (ctx); > >>>>> + /* For functions peel the context up to namespace/TU scope. The > >> abstract > >>>>> + copies reveal the true nesting. */ > >>>>> + if (TREE_CODE (decl) == FUNCTION_DECL) > >>>>> + while (ctx && TREE_CODE (ctx) == FUNCTION_DECL) > >>>>> + ctx = DECL_CONTEXT (ctx); > >>>>> /* Likewise namespaces in case we do not want to emit DIEs for > >> them. */ > >>>>> if (debug_info_level <= DINFO_LEVEL_TERSE) > >>>>> while (ctx && TREE_CODE (ctx) == NAMESPACE_DECL) > >>>>> @@ -6099,8 +6104,7 @@ maybe_create_die_with_external_ref (tree > >> decl) > >>>>> /* Leave function local entities parent determination to > >> when > >>>>> we process scope vars. */ > >>>>> ; > >>>>> - else > >>>>> - parent = lookup_decl_die (ctx); > >>>>> + parent = lookup_decl_die (ctx); > >>>>> } > >>>>> else > >>>>> /* In some cases the FEs fail to set DECL_CONTEXT properly. > >>>>> > >>>> > >>>> Thanks. This solves the problem. Only the first hunk was required. > >> Second hunk > >>>> actually causes an ICE when TREE_CODE (ctx) == BLOCK. > >>>> OK to commit the attached patch? > >>> > >>> I think we need to merge the TYPE_P peeling and FUNCTION_DECL peeling > >> into > >>> one loop since I suppose we can have a nested function in class > >> scope. > >>> So sth like > >>> > >>> diff --git a/gcc/dwarf2out.c b/gcc/dwarf2out.c > >>> index 82783c4968b..61228410b51 100644 > >>> --- a/gcc/dwarf2out.c > >>> +++ b/gcc/dwarf2out.c > >>> @@ -6073,8 +6073,12 @@ maybe_create_die_with_external_ref (tree decl) > >>> } > >>> else > >>> ctx = DECL_CONTEXT (decl); > >>> - /* Peel types in the context stack. */ > >>> - while (ctx && TYPE_P (ctx)) > >>> + /* Peel types in the context stack. For functions peel the > >> context up > >>> + to namespace/TU scope. The abstract copies reveal the true > >> nesting. */ > >>> + while (ctx > >>> + && (TYPE_P (ctx) > >>> + || (TREE_CODE (decl) == FUNCTION_DECL > >>> + && TREE_CODE (ctx) == FUNCTION_DECL))) > >>> ctx = TYPE_CONTEXT (ctx); > >>> /* Likewise namespaces in case we do not want to emit DIEs for > >> them. */ > >>> if (debug_info_level <= DINFO_LEVEL_TERSE) > >>> > >> This causes an ICE, > >> internal compiler error: tree check: expected class 'type', have > >> 'declaration' (function_decl) > >> > >> Did you intend something like this: > >> > >> diff --git a/gcc/dwarf2out.c b/gcc/dwarf2out.c > >> index 561f8b23517..c61f0041fba 100644 > >> --- a/gcc/dwarf2out.c > >> +++ b/gcc/dwarf2out.c > >> @@ -6121,3 +6121,8 @@ maybe_create_die_with_external_ref (tree decl) > >> - /* Peel types in the context stack. */ > >> - while (ctx && TYPE_P (ctx)) > >> - ctx = TYPE_CONTEXT (ctx); > >> + /* Peel types in the context stack. For functions peel the context > >> up > >> + to namespace/TU scope. The abstract copies reveal the true > >> nesting. */ > >> + while (ctx > >> + && (TYPE_P (ctx) > >> + || (TREE_CODE (decl) == FUNCTION_DECL > >> + && TREE_CODE (ctx) == FUNCTION_DECL))) > >> + ctx = TYPE_P (ctx) ? TYPE_CONTEXT (ctx) : DECL_CONTEXT (ctx); > >> + > > > > Yes, of course. > > > >> > >>> if that works it's OK. Can you run it on the gdb testsuite with > >> -flto added > >>> as well please (you need to do before/after comparison since IIRC > >> adding > >>> -flto will add a few fails). > > GDB testsuite shows some extra fails which mainly happen because testcase assumes that you can > access the local variable of enclosing function from the nested function (or omp parallel region). > After this change, the nested functions are no longer children of the enclosing function so those > tests fail. I think you should consult with gdb folks on this - the functions are still children of the enclosing function as seen in the abstract instance. Just the concrete instance is put in another place. But yes, that was what I expected as bad side-effect of the change. Now I wonder how to fix that - even for offloading a "good" debugger could allow debugging both the host and the target and DTRT when printing a variable from the containing function on the target (lookup the variable on the host). So I think we need to get to an agreement between the debug info producer and consumer here. Usually the DWARF spec is not of much help here. Richard. > The problem that prompted this patch happened for parent function that did not have a code range i.e > a notional parent. I was wondering if we should update the ctx only for such parents instead of all > function as we did above. > > Thanks, > -- > Hafiz Abid Qadeer > Mentor, a Siemens Business
On Thu, Jul 22, 2021 at 01:43:49PM +0200, Richard Biener wrote: > So I think we need to get to an agreement between the debug info > producer and consumer here. > Usually the DWARF spec is not of much help here. It is something that needs to be discussed for DWARF 6, currently indeed can be solved only with some DWARF extensions we'd need to invent. Jakub
On Thu, Jul 22, 2021 at 1:48 PM Jakub Jelinek <jakub@redhat.com> wrote: > > On Thu, Jul 22, 2021 at 01:43:49PM +0200, Richard Biener wrote: > > So I think we need to get to an agreement between the debug info > > producer and consumer here. > > Usually the DWARF spec is not of much help here. > > It is something that needs to be discussed for DWARF 6, currently indeed can > be solved only with some DWARF extensions we'd need to invent. I mean, the question is what should the concrete instance inherit from the abstract instance - IMHO parent-child relationship is one thing, no? > Jakub >
On 22/07/2021 12:52, Richard Biener wrote: > On Thu, Jul 22, 2021 at 1:48 PM Jakub Jelinek <jakub@redhat.com> wrote: >> >> On Thu, Jul 22, 2021 at 01:43:49PM +0200, Richard Biener wrote: >>> So I think we need to get to an agreement between the debug info >>> producer and consumer here. >>> Usually the DWARF spec is not of much help here. >> >> It is something that needs to be discussed for DWARF 6, currently indeed can >> be solved only with some DWARF extensions we'd need to invent. > > I mean, the question is what should the concrete instance inherit from > the abstract instance - IMHO parent-child relationship is one thing, no? I guess the problem is that pointer is one-sided from concrete to abstract. With this change, one can go from concrete child function to abstract child (and abstract parent). But it is not easy to find the concrete parent for the consumer as there is no link from abstract to concrete. Thanks,
On Mon, Jul 26, 2021 at 11:34 PM Hafiz Abid Qadeer <abid_qadeer@mentor.com> wrote: > > On 22/07/2021 12:52, Richard Biener wrote: > > On Thu, Jul 22, 2021 at 1:48 PM Jakub Jelinek <jakub@redhat.com> wrote: > >> > >> On Thu, Jul 22, 2021 at 01:43:49PM +0200, Richard Biener wrote: > >>> So I think we need to get to an agreement between the debug info > >>> producer and consumer here. > >>> Usually the DWARF spec is not of much help here. > >> > >> It is something that needs to be discussed for DWARF 6, currently indeed can > >> be solved only with some DWARF extensions we'd need to invent. > > > > I mean, the question is what should the concrete instance inherit from > > the abstract instance - IMHO parent-child relationship is one thing, no? > > I guess the problem is that pointer is one-sided from concrete to abstract. With this change, one > can go from concrete child function to abstract child (and abstract parent). But it is not easy to > find the concrete parent for the consumer as there is no link from abstract to concrete. Yes, that's true - there could be a one-to-many relationship there. But then I wonder in which case such lookup in the DIE tree would be the correct thing to do. If I lookup a variable from the parent then the concrete instance of that should be found by unwinding unless it is a static variable in which case the lookup can be done in any of the concrete instances. But then the original issue that the consumer skips the function if it doesn't have a PC range and thus skips over childs looks like an invalid optimization. Btw, the situation you run into can be simulated by int main(int argc, char **argv) { void foo () { __builtin_puts ("bar"); } foo (); return 0; } and compiling with -g -flto -flto-partition=max which forces main and foo into different LTRANS units and get's us <1><114>: Abbrev Number: 2 (DW_TAG_subprogram) <115> DW_AT_abstract_origin: <0x155> <2><119>: Abbrev Number: 3 (DW_TAG_subprogram) <11a> DW_AT_abstract_origin: <0x179> <11e> DW_AT_low_pc : 0x400515 <126> DW_AT_high_pc : 0x19 <12e> DW_AT_frame_base : 1 byte block: 9c (DW_OP_call_frame_cfa) <130> DW_AT_static_link : 4 byte block: 91 68 6 6 (DW_OP_fbreg: -2 gdb then fails to see 'foo' at all (cannot break on it) and the lookup of 'argc' inside it fails (setting a breakpoint also fails without -flto-partition=max, but I can print argc just fine). I suggest you file a bug with gdb and see how that resolves. Richard. > > Thanks, > -- > Hafiz Abid Qadeer > Mentor, a Siemens Business
On 27/07/2021 09:39, Richard Biener wrote: > Yes, that's true - there could be a one-to-many relationship there. But then I > wonder in which case such lookup in the DIE tree would be the correct thing > to do. If I lookup a variable from the parent then the concrete > instance of that > should be found by unwinding unless it is a static variable in which case > the lookup can be done in any of the concrete instances. > > But then the original issue that the consumer skips the function if it doesn't > have a PC range and thus skips over childs looks like an invalid optimization. > > Btw, the situation you run into can be simulated by > > int main(int argc, char **argv) > { > void foo () > { > __builtin_puts ("bar"); > } > foo (); > return 0; > } > > and compiling with -g -flto -flto-partition=max which forces main and foo > into different LTRANS units and get's us > > <1><114>: Abbrev Number: 2 (DW_TAG_subprogram) > <115> DW_AT_abstract_origin: <0x155> > <2><119>: Abbrev Number: 3 (DW_TAG_subprogram) > <11a> DW_AT_abstract_origin: <0x179> > <11e> DW_AT_low_pc : 0x400515 > <126> DW_AT_high_pc : 0x19 > <12e> DW_AT_frame_base : 1 byte block: 9c (DW_OP_call_frame_cfa) > <130> DW_AT_static_link : 4 byte block: 91 68 6 6 (DW_OP_fbreg: -2 > > gdb then fails to see 'foo' at all (cannot break on it) and the lookup of 'argc' > inside it fails (setting a breakpoint also fails without -flto-partition=max, > but I can print argc just fine). > > I suggest you file a bug with gdb and see how that resolves. I have filed https://sourceware.org/bugzilla/show_bug.cgi?id=28147 for this issue. Thanks,
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. These + functions exist on the host side but not in the offloed code. But they + 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<dwarf_file_data> { typedef const char *compare_type; @@ -23652,8 +23657,23 @@ gen_subprogram_die (tree decl, dw_die_ref context_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 function + 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 the + parent functions here and will later fix them. */ + if (lookup_attribute ("omp target entrypoint", + DECL_ATTRIBUTES (decl))) + { + limbo_die_node *node = ggc_cleared_alloc<limbo_die_node> (); + node->die = subr_die->die_parent; + node->created_for = decl; + node->next = notional_parents_list; + notional_parents_list = node; + } } else { @@ -31881,6 +31901,46 @@ flush_limbo_die_list (void) } } +/* Fixup notional parent function (which does not actually exist) so that + a function with no address range is not parent of a function *with* address + ranges. Otherwise debugger see the parent function without code range + 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 parent + function only exists in the host-side portion of the code. */ + +static void +fixup_notional_parents (void) +{ + limbo_die_node *node; + + for (node = notional_parents_list; node; node = node->next) + { + dw_die_ref notional_parent = node->die; + /* The dwarf at this moment looks like this + DW_TAG_compile_unit + DW_AT_name ("<artificial>") + + 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 == DW_TAG_subprogram + && !(get_AT (notional_parent, DW_AT_low_pc) + || get_AT (notional_parent, DW_AT_ranges))) + + { + dw_die_ref cu = notional_parent->die_parent; + if (cu && cu->die_tag == 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 = NULL; comdat_type_list = NULL; limbo_die_list = NULL; + notional_parents_list = NULL; file_table = NULL; decl_die_table = NULL; common_block_die_table = NULL;
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 ("<artificial>") 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 which 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 in 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 offload function. This results in a poor debug experience of offload code. This patch tries to solve this problem by making offload kernels children 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. 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(-)