From patchwork Thu Jul 1 15:16:57 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Abid Qadeer X-Patchwork-Id: 1499624 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=2620:52:3:1:0:246e:9693:128c; helo=sourceware.org; envelope-from=gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Received: from sourceware.org (server2.sourceware.org [IPv6:2620:52:3:1:0:246e:9693:128c]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature RSA-PSS (4096 bits) server-digest SHA256) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 4GG1zS25SWz9sV8 for ; Fri, 2 Jul 2021 01:17:24 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 733EF396EC6F for ; Thu, 1 Jul 2021 15:17:21 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa2.mentor.iphmx.com (esa2.mentor.iphmx.com [68.232.141.98]) by sourceware.org (Postfix) with ESMTPS id 2CEDA383D818 for ; Thu, 1 Jul 2021 15:17:09 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 2CEDA383D818 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com IronPort-SDR: KkJv8mshcsFpdZGfv0aMx2Qk423cNo4RAfW/VbXcC0ckJH+Ia84cNBXlqGtdYD1fi0Tzr85Xhj FSVKBbcf492FUdQDT4Fw3eo86+mFxEKsGdBLZDQIyVa/5qBL4rPeycQ5FgUuNZTkJ7N0qbTxaS sYCjxTJg6u0wXLYkwnzEe+ZZVfvD5jQ19obtDCu2UbwE9r+kColBaXsIzQ+WGi7iGPauMdWL+v UU1nngzXYNEZotOMG6Ddy2za8DWSMal6W2Ie0KsLBJtWr1/6Xg9JJvK+OETE8wJze/IC772CoW vFk= X-IronPort-AV: E=Sophos;i="5.83,314,1616486400"; d="scan'208";a="63064853" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa2.mentor.iphmx.com with ESMTP; 01 Jul 2021 07:17:07 -0800 IronPort-SDR: ZVOwl2bYf0NgFGD+TPxU2I+OH2o4hngr809ro+KTKpUY4yO7S/u7GiKykHVQJOLl7BGRvWx/A/ TTkOd1sAZ4v6kICkTMjfL53DSaZBFIwtXSUTq5cddE5lEwfrCe7XOwvN2ZpgbZNzmcNzuEIiUx T1zXMki09spN2ePyQrkRWtnUn0uVKU42cphrkOl9fr28b+EkymhLmDSpzTrKsKkPkqEc+dT5Aa fXLJiZTeFclJdAOUf1HOh6VWnZ0q4RbXZ7QrcNZD29O3IjFVBx5AIB77462olj/UVicONMfY28 4B8= From: Hafiz Abid Qadeer To: Subject: [PATCH] [DWARF] Fix hierarchy of debug information for offload kernels. Date: Thu, 1 Jul 2021 16:16:57 +0100 Message-ID: <20210701151657.935006-1-abidh@codesourcery.com> X-Mailer: git-send-email 2.25.1 MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: SVR-IES-MBX-07.mgc.mentorg.com (139.181.222.7) To SVR-IES-MBX-03.mgc.mentorg.com (139.181.222.3) X-Spam-Status: No, score=-12.9 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.4 X-Spam-Checker-Version: SpamAssassin 3.4.4 (2020-01-24) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Cc: abid_qadeer@mentor.com Errors-To: gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org Sender: "Gcc-patches" Currently, if we look at the debug information for offload kernel regions, it looks something like this: void foo (void) { #pragma acc kernels { } } DW_TAG_compile_unit DW_AT_name ("") DW_TAG_subprogram // notional parent function (foo) with no code range DW_TAG_subprogram // offload function foo._omp_fn.0 There is an artificial compile unit. It contains a parent subprogram 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 { 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 (); + 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 ("") + + 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;