[PATCH] [DWARF] Fix hierarchy of debug information for offload kernels.

[PATCH] [DWARF] Fix hierarchy of debug information for offload kernels.

[Hafiz Abid Qadeer]

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(-)

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

Re: [PATCH] [DWARF] Fix hierarchy of debug information for offload kernels.

[Richard Biener]

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
>

Re: [PATCH] [DWARF] Fix hierarchy of debug information for offload kernels.

[Thomas Schwinge]

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

Re: [PATCH] [DWARF] Fix hierarchy of debug information for offload kernels.

[Hafiz Abid Qadeer]

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.

-- 
Hafiz Abid Qadeer
Mentor, a Siemens Business

Re: [PATCH] [DWARF] Fix hierarchy of debug information for offload kernels.

[Richard Biener]

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

Re: [PATCH] [DWARF] Fix hierarchy of debug information for offload kernels.

[Hafiz Abid Qadeer]

[-- Attachment #1: Type: text/plain, Size: 2449 bytes --]

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?


-- 
Hafiz Abid Qadeer
Mentor, a Siemens Business

[-- Attachment #2: 0001-DWARF-Fix-hierarchy-of-debug-information-for-offload.patch --]
[-- Type: text/x-patch, Size: 2134 bytes --]

From 8e886f8502784d3aafdaf7e9778ce21b8c8f3b93 Mon Sep 17 00:00:00 2001
From: Hafiz Abid Qadeer <abidh@codesourcery.com>
Date: Fri, 16 Jul 2021 21:00:37 +0100
Subject: [PATCH] [DWARF] Fix hierarchy of debug information for offload kernels.

Currently, if we look at the debug information for offload kernel
regions, it looks something like this:

void foo (void)
{
  {

  }
}

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 was suggested by Richard and it solves this problem by peeling
the parent function from the concrete copies.

gcc/

	* gcc/dwarf2out.c (maybe_create_die_with_external_ref): Remove function
	from the context chain.
---
 gcc/dwarf2out.c | 5 +++++
 1 file changed, 5 insertions(+)

diff --git a/gcc/dwarf2out.c b/gcc/dwarf2out.c
index 561f8b23517..e2893bd91ed 100644
--- a/gcc/dwarf2out.c
+++ b/gcc/dwarf2out.c
@@ -6121,6 +6121,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)
-- 
2.25.1

Re: [PATCH] [DWARF] Fix hierarchy of debug information for offload kernels.

[Richard Biener]

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.

Re: [PATCH] [DWARF] Fix hierarchy of debug information for offload kernels.

[Hafiz Abid Qadeer]

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.
> 


-- 
Hafiz Abid Qadeer
Mentor, a Siemens Business

Re: [PATCH] [DWARF] Fix hierarchy of debug information for offload kernels.

[Richard Biener]

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.
>> 

Re: [PATCH] [DWARF] Fix hierarchy of debug information for offload kernels.

[Hafiz Abid Qadeer]

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,
-- 
Hafiz Abid Qadeer
Mentor, a Siemens Business

Re: [PATCH] [DWARF] Fix hierarchy of debug information for offload kernels.

[Richard Biener]

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

Re: [PATCH] [DWARF] Fix hierarchy of debug information for offload kernels.

[Jakub Jelinek]

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

Re: [PATCH] [DWARF] Fix hierarchy of debug information for offload kernels.

[Richard Biener]

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
>

Re: [PATCH] [DWARF] Fix hierarchy of debug information for offload kernels.

[Hafiz Abid Qadeer]

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,
-- 
Hafiz Abid Qadeer
Mentor, a Siemens Business

Re: [PATCH] [DWARF] Fix hierarchy of debug information for offload kernels.

[Richard Biener]

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

Re: [PATCH] [DWARF] Fix hierarchy of debug information for offload kernels.

[Hafiz Abid Qadeer]

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,
-- 
Hafiz Abid Qadeer
Mentor, a Siemens Business