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

[Richard Biener <richard.guenther@gmail.com>]

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
>