Hi!
On 2021-07-02T09:15:27+0200, Richard Biener via Gcc-patches
<[email protected]> wrote:
> On Thu, Jul 1, 2021 at 5:17 PM Hafiz Abid Qadeer <[email protected]>
> 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
