From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from simark.ca by simark.ca with LMTP id 9qifLEsgtWotyzwAWB0awg (envelope-from ) for ; Thu, 24 Sep 2026 09:06:19 -0400 Authentication-Results: simark.ca; dkim=pass (1024-bit key; unprotected) header.d=amd.com header.i=@amd.com header.a=rsa-sha256 header.s=selector1 header.b=fkpcsNRK; dkim-atps=neutral Received: by simark.ca (Postfix, from userid 112) id A0C1B1E01F; Thu, 24 Sep 2026 09:06:19 -0400 (EDT) X-Spam-Checker-Version: SpamAssassin 4.0.1 (2024-03-25) on simark.ca X-Spam-Level: X-Spam-Status: No, score=-6.4 required=5.0 tests=ARC_SIGNED,ARC_VALID,BAYES_00, DKIMWL_WL_HIGH,DKIM_SIGNED,DKIM_VALID,DKIM_VALID_AU,MAILING_LIST_MULTI, RCVD_IN_DNSWL_MED autolearn=ham autolearn_force=no version=4.0.1 Received: from vm01.sourceware.org (vm01.sourceware.org [IPv6:2620:52:6:3111::32]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange x25519 server-signature ECDSA (prime256v1) server-digest SHA256) (No client certificate requested) by simark.ca (Postfix) with ESMTPS id A6F6E1E01F for ; Thu, 24 Sep 2026 09:06:17 -0400 (EDT) Received: from vm01.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 56DC64BB5899 for ; Thu, 24 Sep 2026 13:06:09 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 56DC64BB5899 Authentication-Results: sourceware.org; dkim=pass (1024-bit key, unprotected) header.d=amd.com header.i=@amd.com header.a=rsa-sha256 header.s=selector1 header.b=fkpcsNRK Received: from CY3PR05CU001.outbound.protection.outlook.com (mail-westcentralusazlp170130007.outbound.protection.outlook.com [IPv6:2a01:111:f403:c112::7]) by sourceware.org (Postfix) with ESMTPS id 1AE594BB3BE3 for ; Thu, 24 Sep 2026 13:05:22 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 1AE594BB3BE3 Authentication-Results: sourceware.org; dmarc=pass (p=quarantine dis=none) header.from=amd.com Authentication-Results: sourceware.org; spf=fail smtp.mailfrom=amd.com ARC-Filter: OpenARC Filter v1.0.0 sourceware.org 1AE594BB3BE3 Authentication-Results: sourceware.org; arc=pass smtp.remote-ip=2a01:111:f403:c112::7 ARC-Seal: i=2; a=rsa-sha256; d=sourceware.org; s=key; t=1790255128; cv=pass; b=hZR+SP+HbzW0bLVoMvrNamtCpuBkz6Ffpui0XiMthXhXtg3gfu0Ff/z30k/4JDaCvxvW6i284iP68VxMa7RvjQxiMBn6/TgyVv/wuxYGiQQS1ATS6EK/9cUpQwNbmTgLtW/5pCoCSCcZDQ1mDAFfEdr5THyDjnwP6D+kTehrc1o= ARC-Message-Signature: i=2; a=rsa-sha256; d=sourceware.org; s=key; t=1790255128; c=relaxed/simple; bh=LqSM0EkGkrUvlNmi4405u6nejjsgKEOquZQu0BJ/XYQ=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=qkn1kqWNnbMOvh/1wblSbaxWmNLpWZE5ggbt1/egjgxJVBE18OG07toSwtZhf5qPrMLbhDA6ubCZfBOzYU1JQ6J3cMKxe4CGbxK4effbdagLjeC8XB3n+s9VcJfzFxoPxjuTt6GDH/uDVg1POaBQrzYR660BNU3hW8l6E8stKYI= ARC-Authentication-Results: i=2; sourceware.org; dkim=pass (1024-bit key, unprotected) header.d=amd.com header.i=@amd.com header.a=rsa-sha256 header.s=selector1 header.b=fkpcsNRK DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 1AE594BB3BE3 ARC-Seal: i=1; a=rsa-sha256; s=arcselector10001; d=microsoft.com; cv=none; b=Z3hwq2Ye3NJMgDNhjyFn+oOcvX9cuw0lxNA1bs5llLYkJTcuvBbKC5Ht2Zu8kYs7h9iHX/ozyASlJJHxADchIGYKeKsb0wMDPgPU+l4qGDfLNhyG7RsKliujg/jSvOMKi8DxjHTFbfvmyptIGwT7Zhzb9xxPLN/a/l86H6UFC7w0XetXyZYlQrXhUfVBuCrZ6GUbE0PhKjflc/a+gatCE8P0xPvM8QC5iCLz1LS9SZUJ/2eBCFSahArKkZhGbH1ZoSEcYNZDgw1W40aHwRleaafETWWi9L1S71PygwimHLBwJ01OL2ujC+K+DomIFyXuAKWxYGgGvsWDTyTdzzXzAQ== ARC-Message-Signature: i=1; a=rsa-sha256; c=relaxed/relaxed; d=microsoft.com; s=arcselector10001; h=From:Date:Subject:Message-ID:Content-Type:MIME-Version:X-MS-Exchange-AntiSpam-MessageData-ChunkCount:X-MS-Exchange-AntiSpam-MessageData-0:X-MS-Exchange-AntiSpam-MessageData-1; bh=DiwrpAl0pi6JqenOXEKPvx2Diu+7l5BOa8Y+YdpqeZQ=; b=mtudhAr3m2sqjU4dnmiRvY34mTfXV+RorEWf1YrPb3+271oFkPtlc7lM1gao6TK+n4koa7szEAi4nmEdyq0+j1ChR4yl7dBqQkR00V66U19L1AMhaU5MPW+VvnOkYb0LWTC+VAo0rjHtgyqJkLXwFx7QElwNuJ0qSHhL0jadM2sPJ3p8DzC0HJhSVkPUXGfQ9cbn+/4LtvH0JAU5dmJcrr/qaW5+ejnqwQZi7QOAiFJg4wVxxrPT2mlWcVwNU3Z/jiaIpgNWzJT7hEbbixlnzacpHalcqbOfNcVSnurhjJ+CTT5+fd4TpbFmuap3uJo7nvn4L3C6sCHwuJ/8+BJovw== ARC-Authentication-Results: i=1; mx.microsoft.com 1; spf=pass (sender ip is 165.204.84.17) smtp.rcpttodomain=sourceware.org smtp.mailfrom=amd.com; dmarc=pass (p=quarantine sp=quarantine pct=100) action=none header.from=amd.com; dkim=none (message not signed); arc=none (0) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=amd.com; s=selector1; h=From:Date:Subject:Message-ID:Content-Type:MIME-Version:X-MS-Exchange-SenderADCheck; bh=DiwrpAl0pi6JqenOXEKPvx2Diu+7l5BOa8Y+YdpqeZQ=; b=fkpcsNRKPdY2DEwq4SYtam3PkcLvWiw/Eav9YCRRib9JlPDs6szjEZONQTElaIM4aS6HMAWl1K4xghQyiSu9931zQ6/m0ph6E705KF+QPFgvdfWv86+VQux5Vf8HqLrRQbQ8EiOSv+33rTFGSh6mMB7ZfmDnc+sENoKbOwYImrE= Received: from MN0PR04CA0011.namprd04.prod.outlook.com (2603:10b6:208:52d::7) by SA0PR12MB4432.namprd12.prod.outlook.com (2603:10b6:806:98::16) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_256_GCM_SHA384) id 15.21.451.18; Thu, 24 Sep 2026 13:05:15 +0000 Received: from BN3PEPF0000B074.namprd04.prod.outlook.com (2603:10b6:208:52d:cafe::32) by MN0PR04CA0011.outlook.office365.com (2603:10b6:208:52d::7) with Microsoft SMTP Server (version=TLS1_3, cipher=TLS_AES_256_GCM_SHA384) id 15.21.451.16 via Frontend Transport; Thu, 24 Sep 2026 13:05:14 +0000 X-MS-Exchange-Authentication-Results: mx.microsoft.com 1; spf=pass (sender IP is 165.204.84.17) smtp.mailfrom=amd.com; dkim=none (message not signed) header.d=none;dmarc=pass action=none header.from=amd.com; Received-SPF: Pass (protection.outlook.com: domain of amd.com designates 165.204.84.17 as permitted sender) receiver=protection.outlook.com; client-ip=165.204.84.17; helo=satlexmb07.amd.com; pr=C Received: from satlexmb07.amd.com (165.204.84.17) by BN3PEPF0000B074.mail.protection.outlook.com (10.167.243.119) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_256_GCM_SHA384) id 15.21.451.8 via Frontend Transport; Thu, 24 Sep 2026 13:05:14 +0000 Received: from RSBBFILIPOV02.amd.com (10.180.168.240) by satlexmb07.amd.com (10.181.42.216) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_256_GCM_SHA384) id 15.2.2562.49; Thu, 24 Sep 2026 08:05:12 -0500 From: Bratislav Filipovic To: CC: , , , , Bratislav Filipovic Subject: [PATCH] gdb: Implement stop-on-solib-events for GPU code objects Date: Thu, 24 Sep 2026 15:04:51 +0200 Message-ID: <20260924130451.89610-1-bfilipov@amd.com> X-Mailer: git-send-email 2.43.0 MIME-Version: 1.0 Content-Transfer-Encoding: 8bit Content-Type: text/plain X-Originating-IP: [10.180.168.240] X-ClientProxiedBy: satlexmb07.amd.com (10.181.42.216) To satlexmb07.amd.com (10.181.42.216) X-EOPAttributedMessage: 0 X-MS-PublicTrafficType: Email X-MS-TrafficTypeDiagnostic: BN3PEPF0000B074:EE_|SA0PR12MB4432:EE_ X-MS-Office365-Filtering-Correlation-Id: 88b317d2-7d09-40ed-421c-08df1a3c7987 X-MS-Exchange-SenderADCheck: 1 X-MS-Exchange-AntiSpam-Relay: 0 X-Microsoft-Antispam: BCL:0; ARA:13230040|82310400026|376014|23010399003|1800799024|36860700016|6133799003|13003099007|18002099003|11063799006|5023799004|56012099006|10067099003|3023799007; X-Microsoft-Antispam-Message-Info: MZEol88z6SLdL+AJeoYZ3VK6/GDszlTvFwNxO7VS0DTovvj6vS0WLU/EhoQ7VxQy+alR1LGx3YGg4gvQ7dy67JNjFPQqUxDgxX4Z5jpqBs6uRxm7ARafSfIdFn+WtLJAHZZHkqLnpAhY7zqunLVur/LVDJuByd2Fyop4zR7zuCXVYR/Gt6/HwopBtE5MP3EuO7ZQiDNmidWTHFxAvnYEe38NGW+dw1ZZ9f+kaAvyyAG7l9gZ1M7ZA+D/3y4yauZkgMSBFNfExmBniklVB+o0tdBlE20gTRaE9uwCtzUldG+8H9XiPu72jzRMAG8ocGmd//5mbI/eVyEwdgsId+RFGB4PAbU2WGdzDOxhzq7PvrM0GirfT12badVFC711/mq7KZ/vxSqupzKseWfiJZleJYOHeDsz3NQlekt9qnYVMBilcNsogZ+CDJ1wOI3aF/Kn4cHAsScOtH69wmkoo8CNJawSe7Lx47MqFA3HnHyiLHnoBx5pe7NU31qhpL1ZyHATq3QJMVEAU6zUhgT+9i29uEIsy4hwcZFwrrfe0ok32DAU4EXXrlhsgSwjAVbZGprqUEhia37DiEDKskcuNchyEYZmMD8AWTbtVFFUPrt0J79wi29PsNRJNW0In57UH6OtbvkoHDvxfr21bboVBIkJLg== X-Forefront-Antispam-Report: CIP:165.204.84.17; CTRY:US; LANG:en; SCL:1; SRV:; IPV:NLI; SFV:NSPM; H:satlexmb07.amd.com; PTR:InfoDomainNonexistent; CAT:NONE; SFS:(13230040)(82310400026)(376014)(23010399003)(1800799024)(36860700016)(6133799003)(13003099007)(18002099003)(11063799006)(5023799004)(56012099006)(10067099003)(3023799007); DIR:OUT; SFP:1101; X-MS-Exchange-AntiSpam-MessageData-ChunkCount: 1 X-MS-Exchange-AntiSpam-MessageData-0: hxMTajvwzpv9BHiSSNj1K7tOBnsLn7LRg+7dcKggczAyMiZCIxj7Rd0fSvFGEgp3OBnzE0JDW5e+nYUHGN2G2PcGP+q3osbi7+RhV0Osz9ECPs3qryKBIL+WMYUSa219RtTqqCTwr0jWkKusiegED1ggcL0sIREYiXjYvTfygp2KLcNU4sSjVkHNx9a70zZXHdeKZqNa3dJ+FuCh2j4kKnEQgS85MgrgCw1gcrIwqv+6qvBDxnUd0ivBQmq8cmt18VCs1StJZEjwbJv06hk3T/ib+pJXMYLWdLZ30Bewk9k5H+oJnMUjO0FITOoaQepwpSsvdO9WANnNnHVAOJhNlY7JPZ5YsJoYhycvhlkvU1ylWJevaPL/79VCXEi9O1VMpmPbIKUYmXceRSbTf/hqGpDexZmc5olY/ATCfX/Ddd8SWoqcTUYmcxzQXz0yTmFU X-OriginatorOrg: amd.com X-MS-Exchange-CrossTenant-OriginalArrivalTime: 24 Sep 2026 13:05:14.8729 (UTC) X-MS-Exchange-CrossTenant-Network-Message-Id: 88b317d2-7d09-40ed-421c-08df1a3c7987 X-MS-Exchange-CrossTenant-Id: 3dd8961f-e488-4e60-8e11-a82d994e183d X-MS-Exchange-CrossTenant-OriginalAttributedTenantConnectingIp: TenantId=3dd8961f-e488-4e60-8e11-a82d994e183d; Ip=[165.204.84.17]; Helo=[satlexmb07.amd.com] X-MS-Exchange-CrossTenant-AuthSource: BN3PEPF0000B074.namprd04.prod.outlook.com X-MS-Exchange-CrossTenant-AuthAs: Anonymous X-MS-Exchange-CrossTenant-FromEntityHeader: HybridOnPrem X-MS-Exchange-Transport-CrossTenantHeadersStamped: SA0PR12MB4432 X-BeenThere: gdb-patches@sourceware.org X-Mailman-Version: 2.1.30 Precedence: list List-Id: Gdb-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gdb-patches-bounces~public-inbox=simark.ca@sourceware.org GDB's "set stop-on-solib-events 1" setting allows users to stop execution when shared libraries are loaded or unloaded, enabling inspection and breakpoint placement before library code executes. This feature works for CPU shared libraries but is not working for GPU code objects loaded by the AMD ROCm runtime. This commit implements stop-on-solib-events support for GPU code objects, making GPU code object load/unload events behave consistently with CPU shared library events. The root cause is in amd_dbgapi_target_breakpoint::check_status(), which unconditionally sets bs->stop = 0 and bs->print_it = print_it_noop, regardless of the stop_on_solib_events setting. This is in contrast to internal_breakpoint::check_status() for CPU shared libraries, which respects the setting. The fix includes: 1. Generalize print_solib_event() function to eliminate code duplication between CPU and GPU event printing. The function now accepts parameters for event description, field names, and plural forms. This reduces ~50 lines of duplicated code. 2. Add code_object_list_updated flag to amd_dbgapi_inferior_info to track when AMD_DBGAPI_EVENT_KIND_CODE_OBJECT_LIST_UPDATED events occur during process_event_queue(). 3. Modify check_status() to check this flag after processing events and update bs->stop, bs->print, and bs->print_it based on stop_on_solib_events. 4. Implement print_it() override to display "Stopped due to GPU code object event" to distinguish GPU events from CPU shared library events. This makes GPU code object load events behave consistently with CPU shared library events. A test is included in gdb.rocm/solib-event.exp. --- gdb/NEWS | 6 + gdb/amd-dbgapi-target.c | 83 +++++++++++++ gdb/breakpoint.c | 31 +++-- gdb/breakpoint.h | 6 +- gdb/doc/gdb.texinfo | 9 +- gdb/infrun.c | 5 +- gdb/testsuite/gdb.rocm/solib-event-kernel.cpp | 27 +++++ gdb/testsuite/gdb.rocm/solib-event.cpp | 105 ++++++++++++++++ gdb/testsuite/gdb.rocm/solib-event.exp | 112 ++++++++++++++++++ 9 files changed, 369 insertions(+), 15 deletions(-) create mode 100644 gdb/testsuite/gdb.rocm/solib-event-kernel.cpp create mode 100644 gdb/testsuite/gdb.rocm/solib-event.cpp create mode 100644 gdb/testsuite/gdb.rocm/solib-event.exp diff --git a/gdb/NEWS b/gdb/NEWS index 77ad2d3fc22..957ba2e8563 100644 --- a/gdb/NEWS +++ b/gdb/NEWS @@ -3,6 +3,12 @@ *** Changes since GDB 18 +* The 'set stop-on-solib-events' setting now applies to GPU code objects + loaded by the AMD ROCm runtime, in addition to CPU shared libraries. + When enabled (non-zero value), GDB will stop execution whenever GPU code + objects are loaded or unloaded, allowing inspection and breakpoint + placement. + * GDB now distinguishes between the GNU (MinGW) and MSVC Windows ABIs. The "set osabi" command accepts two new values, "Windows-GNU" and diff --git a/gdb/amd-dbgapi-target.c b/gdb/amd-dbgapi-target.c index 93207a8fdd2..0380c4c88d5 100644 --- a/gdb/amd-dbgapi-target.c +++ b/gdb/amd-dbgapi-target.c @@ -30,11 +30,13 @@ #include "gdbsupport/unordered_map.h" #include "inf-loop.h" #include "inferior.h" +#include "infrun.h" #include "objfiles.h" #include "observable.h" #include "registry.h" #include "solib.h" #include "target.h" +#include "ui-out.h" #include @@ -104,6 +106,10 @@ amd_dbgapi_lib_debug_module () scoped_debug_start_end (debug_amd_dbgapi, amd_dbgapi_debug_module (), \ fmt, ##__VA_ARGS__) +/* MI field value for GPU code object kind. */ + +static const char *gpu_code_object_kind = "gpu-code-object"; + /* inferior_created observer token. */ static gdb::observers::token amd_dbgapi_target_inferior_created_observer_token; @@ -258,6 +264,11 @@ struct amd_dbgapi_inferior_info /* List of pending events the amd-dbgapi target retrieved from the dbgapi. */ std::list> wave_events; + /* Flag to track if a CODE_OBJECT_LIST_UPDATED event was seen during + process_event_queue. Used to implement stop-on-solib-events for GPU + code objects. */ + bool code_object_list_updated = false; + /* Map of threads with ongoing displaced steps to corresponding amd-dbgapi displaced stepping handles. */ gdb::unordered_mapstop = true; + bs->print = true; + /* Allow print_it () to print the GPU code object event message. */ + bs->print_it = print_it_normal; + } +} + +enum print_stop_action +amd_dbgapi_target_breakpoint::print_it (const bpstat *bs) const +{ + /* We only reach here when check_status set bs->print_it to print_it_normal, + which happens only for GPU code object events when stop_on_solib_events + is enabled. */ + bool any_deleted = !current_program_space->deleted_solibs.empty (); + bool any_added = !current_program_space->added_solibs.empty (); + + if (any_added || any_deleted) + current_uiout->text (_("Stopped due to GPU code object event:\n")); + else + current_uiout->text (_("Stopped due to GPU code object event (no " + "code objects added or removed)\n")); + + if (current_uiout->is_mi_like_p ()) + { + current_uiout->field_string + ("reason", async_reason_lookup (EXEC_ASYNC_SOLIB_EVENT)); + current_uiout->field_string ("object-kind", gpu_code_object_kind); + } + + if (any_deleted) + { + current_uiout->text (_(" Inferior unloaded ")); + ui_out_emit_list list_emitter (current_uiout, "removed"); + bool first = true; + for (const std::string &name : current_program_space->deleted_solibs) + { + if (!first) + current_uiout->text (" "); + first = false; + current_uiout->field_string ("code-object", name); + current_uiout->text ("\n"); + } + } + + if (any_added) + { + current_uiout->text (_(" Inferior loaded ")); + ui_out_emit_list list_emitter (current_uiout, "added"); + bool first = true; + for (solib *iter : current_program_space->added_solibs) + { + if (!first) + current_uiout->text (" "); + first = false; + current_uiout->field_string ("code-object", iter->name); + current_uiout->text ("\n"); + } + } + + return PRINT_NOTHING; } bool @@ -1565,6 +1647,7 @@ process_one_event (amd_dbgapi_inferior_info &info, inferior is the inferior that hit the breakpoint, which should still be the case now. */ gdb_assert (info.inf == current_inferior ()); + info.code_object_list_updated = true; handle_solib_event (); break; diff --git a/gdb/breakpoint.c b/gdb/breakpoint.c index 988737b6f96..05f55db6a7e 100644 --- a/gdb/breakpoint.c +++ b/gdb/breakpoint.c @@ -5113,7 +5113,9 @@ print_bp_stop_message (bpstat *bs) /* See breakpoint.h. */ void -print_solib_event (bool is_catchpoint) +print_solib_event (bool is_catchpoint, const char *event_description, + const char *plural_description, const char *item_field_name, + const char *object_kind) { bool any_deleted = !current_program_space->deleted_solibs.empty (); bool any_added = !current_program_space->added_solibs.empty (); @@ -5121,15 +5123,28 @@ print_solib_event (bool is_catchpoint) if (!is_catchpoint) { if (any_added || any_deleted) - current_uiout->text (_("Stopped due to shared library event:\n")); + { + std::string msg = string_printf (_("Stopped due to %s event:\n"), + event_description); + current_uiout->text (msg.c_str ()); + } else - current_uiout->text (_("Stopped due to shared library event (no " - "libraries added or removed)\n")); + { + std::string msg = string_printf (_("Stopped due to %s event (no " + "%s added or removed)\n"), + event_description, + plural_description); + current_uiout->text (msg.c_str ()); + } } if (current_uiout->is_mi_like_p ()) - current_uiout->field_string ("reason", - async_reason_lookup (EXEC_ASYNC_SOLIB_EVENT)); + { + current_uiout->field_string + ("reason", async_reason_lookup (EXEC_ASYNC_SOLIB_EVENT)); + if (object_kind != nullptr) + current_uiout->field_string ("object-kind", object_kind); + } if (any_deleted) { @@ -5141,7 +5156,7 @@ print_solib_event (bool is_catchpoint) if (ix > 0) current_uiout->text (" "); - current_uiout->field_string ("library", name); + current_uiout->field_string (item_field_name, name); current_uiout->text ("\n"); } } @@ -5156,7 +5171,7 @@ print_solib_event (bool is_catchpoint) if (!first) current_uiout->text (" "); first = false; - current_uiout->field_string ("library", iter->name); + current_uiout->field_string (item_field_name, iter->name); current_uiout->text ("\n"); } } diff --git a/gdb/breakpoint.h b/gdb/breakpoint.h index 75de448b16e..d7078a0e198 100644 --- a/gdb/breakpoint.h +++ b/gdb/breakpoint.h @@ -2088,7 +2088,11 @@ extern void catch_exception_event (enum exception_event_kind ex_event, IS_CATCHPOINT is true if the event is due to a "catch load" catchpoint, false otherwise. */ -extern void print_solib_event (bool is_catchpoint); +extern void print_solib_event (bool is_catchpoint, + const char *event_description = "shared library", + const char *plural_description = "libraries", + const char *item_field_name = "library", + const char *object_kind = nullptr); /* Print a message describing any user-breakpoints set at PC. This concerns with logical breakpoints, so we match program spaces, not diff --git a/gdb/doc/gdb.texinfo b/gdb/doc/gdb.texinfo index a24f67cb8de..2c948737720 100644 --- a/gdb/doc/gdb.texinfo +++ b/gdb/doc/gdb.texinfo @@ -22513,14 +22513,15 @@ The surrounding square brackets are optional. @item set stop-on-solib-events @kindex set stop-on-solib-events This command controls whether @value{GDBN} should give you control -when the dynamic linker notifies it about some shared library event. -The most common event of interest is loading or unloading of a new -shared library. +when the dynamic linker notifies it about some shared library event, +or when GPU code objects are loaded or unloaded (AMD ROCm targets). +The most common events of interest are loading or unloading of a new +shared library or code object. @item show stop-on-solib-events @kindex show stop-on-solib-events Show whether @value{GDBN} stops and gives you control when shared -library events happen. +library events or GPU code object events happen. @end table Shared libraries are also supported in many cross or remote debugging diff --git a/gdb/infrun.c b/gdb/infrun.c index 92b21017b03..5e00f9d653c 100644 --- a/gdb/infrun.c +++ b/gdb/infrun.c @@ -10849,8 +10849,9 @@ leave it stopped or free to run as needed."), Set stopping for shared library events."), _("\ Show stopping for shared library events."), _("\ If nonzero, gdb will give control to the user when the dynamic linker\n\ -notifies gdb of shared library events. The most common event of interest\n\ -to the user would be loading/unloading of a new library."), +notifies gdb of shared library events, or when GPU code objects are loaded\n\ +or unloaded (AMD ROCm targets). The most common events of interest to the\n\ +user would be loading/unloading of a new library or code object."), set_stop_on_solib_events, show_stop_on_solib_events, &setlist, &showlist); diff --git a/gdb/testsuite/gdb.rocm/solib-event-kernel.cpp b/gdb/testsuite/gdb.rocm/solib-event-kernel.cpp new file mode 100644 index 00000000000..b2c47e26879 --- /dev/null +++ b/gdb/testsuite/gdb.rocm/solib-event-kernel.cpp @@ -0,0 +1,27 @@ +/* This testcase is part of GDB, the GNU debugger. + + Copyright 2026 Free Software Foundation, Inc. + + This program is free software; you can redistribute it and/or modify + it under the terms of the GNU General Public License as published by + the Free Software Foundation; either version 3 of the License, or + (at your option) any later version. + + This program is distributed in the hope that it will be useful, + but WITHOUT ANY WARRANTY; without even the implied warranty of + MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the + GNU General Public License for more details. + + You should have received a copy of the GNU General Public License + along with this program. If not, see . */ + +/* Device kernel for solib-event test. + Compiled with --cuda-device-only to produce .co file. */ + +#include + +extern "C" __global__ void +test_kernel () +{ + asm volatile ("s_nop 1"); +} diff --git a/gdb/testsuite/gdb.rocm/solib-event.cpp b/gdb/testsuite/gdb.rocm/solib-event.cpp new file mode 100644 index 00000000000..823ab5629e1 --- /dev/null +++ b/gdb/testsuite/gdb.rocm/solib-event.cpp @@ -0,0 +1,105 @@ +/* This testcase is part of GDB, the GNU debugger. + + Copyright 2026 Free Software Foundation, Inc. + + This program is free software; you can redistribute it and/or modify + it under the terms of the GNU General Public License as published by + the Free Software Foundation; either version 3 of the License, or + (at your option) any later version. + + This program is distributed in the hope that it will be useful, + but WITHOUT ANY WARRANTY; without even the implied warranty of + MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the + GNU General Public License for more details. + + You should have received a copy of the GNU General Public License + along with this program. If not, see . */ + +#include "rocm-test-utils.h" +#include +#include +#include + +/* Test file:// events (hipModuleLoad). */ +static void +test_file_load (const char *module_path) +{ + hipModule_t module; + CHECK (hipModuleLoad (&module, module_path)); + + hipFunction_t function; + CHECK (hipModuleGetFunction (&function, module, "test_kernel")); + + CHECK (hipModuleLaunchKernel (function, 1, 1, 1, 1, 1, 1, + 0, nullptr, nullptr, nullptr)); + + CHECK (hipDeviceSynchronize ()); + CHECK (hipModuleUnload (module)); +} + +/* Test memory:// events (hipModuleLoadData). */ +static void +test_memory_load (const char *module_path) +{ + /* Read module file into memory buffer. */ + std::ifstream mod (module_path, std::ios::binary | std::ios::ate); + if (!mod.is_open ()) + { + fprintf (stderr, "Failed to open module file\n"); + exit (EXIT_FAILURE); + } + + size_t module_size = mod.tellg (); + mod.seekg (0, std::ios::beg); + std::vector module_buffer (module_size); + + if (!mod.read (reinterpret_cast (module_buffer.data ()), + module_size)) + { + fprintf (stderr, "Failed to read module into memory\n"); + exit (EXIT_FAILURE); + } + mod.close (); + + /* Load from memory buffer. */ + hipModule_t module; + CHECK (hipModuleLoadData (&module, module_buffer.data ())); + + hipFunction_t function; + CHECK (hipModuleGetFunction (&function, module, "test_kernel")); + + CHECK (hipModuleLaunchKernel (function, 1, 1, 1, 1, 1, 1, + 0, nullptr, nullptr, nullptr)); + + CHECK (hipDeviceSynchronize ()); + CHECK (hipModuleUnload (module)); +} + +__global__ void +warm_up () +{ +} + +int +main (int argc, char **argv) +{ + if (argc != 2) + { + fprintf (stderr, "Usage: %s \n", argv[0]); + return EXIT_FAILURE; + } + + const char *module_path = argv[1]; + + /* Submit an empty kernel to force the runtime to do any necessary + setup which might include loading internal code objects. */ + warm_up<<<1, 1>>>(); + CHECK (hipDeviceSynchronize ()); + + /* Enable SOLIB events here. */ + + test_file_load (module_path); + test_memory_load (module_path); + + return 0; +} diff --git a/gdb/testsuite/gdb.rocm/solib-event.exp b/gdb/testsuite/gdb.rocm/solib-event.exp new file mode 100644 index 00000000000..d306cf9f5ff --- /dev/null +++ b/gdb/testsuite/gdb.rocm/solib-event.exp @@ -0,0 +1,112 @@ +# Copyright 2026 Free Software Foundation, Inc. +# Copyright 2026 Advanced Micro Devices, Inc. + +# This program is free software; you can redistribute it and/or modify +# it under the terms of the GNU General Public License as published by +# the Free Software Foundation; either version 3 of the License, or +# (at your option) any later version. +# +# This program is distributed in the hope that it will be useful, +# but WITHOUT ANY WARRANTY; without even the implied warranty of +# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the +# GNU General Public License for more details. +# +# You should have received a copy of the GNU General Public License +# along with this program. If not, see . + +# Test that "set stop-on-solib-events 1" works for GPU code objects. +# Tests both file:// events (hipModuleLoad) and memory:// events +# (hipModuleLoadData), as well as unload events (hipModuleUnload). +# Also tests that the default setting (0) does not stop execution. + +load_lib rocm.exp + +require allow_hip_tests + +standard_testfile .cpp + +# Build device module (.co file). +set kernel_srcfile ${testfile}-kernel.cpp +set hipmodule_path [standard_output_file ${testfile}.co] + +if {[gdb_compile $srcdir/$subdir/$kernel_srcfile \ + $hipmodule_path object \ + {debug hip additional_flags=--cuda-device-only}] != ""} { + return +} + +# Build host executable. +if {[build_executable "failed to prepare" $testfile $srcfile \ + {debug hip}] == -1} { + return +} + +proc test_disabled {} { + with_rocm_gpu_lock { + clean_restart $::testfile + + if {![runto_main -inferior-args $::hipmodule_path]} { + return + } + + # Set breakpoint at program exit and verify execution continues + # through all code object load/unload events without stopping. + gdb_breakpoint [gdb_get_line_number "return 0;"] + gdb_test "continue" \ + "Breakpoint .* main .*return 0.*" \ + "no stops for code object events when disabled" + } +} + +proc test_enabled {} { + with_rocm_gpu_lock { + clean_restart $::testfile + + if {![runto_main -inferior-args $::hipmodule_path]} { + return + } + + gdb_breakpoint [gdb_get_line_number "Enable SOLIB events here"] -temporary + gdb_continue_to_breakpoint "at enable solib-event" + + gdb_test_no_output "set stop-on-solib-events 1" + + with_test_prefix "file://" { + gdb_breakpoint "test_file_load" -temporary + gdb_continue_to_breakpoint "at test_file_load" + + # Test 1: file:// load event (hipModuleLoad). + gdb_test "continue" \ + "Stopped due to GPU code object event.*Inferior loaded file://\[^\r\n\]+" \ + "load" + + # Test 2: file:// unload event (hipModuleUnload). + gdb_test "continue" \ + "Stopped due to GPU code object event.*Inferior unloaded file://\[^\r\n\]+" \ + "unload" + } + + with_test_prefix "memory://" { + gdb_breakpoint "test_memory_load" -temporary + gdb_continue_to_breakpoint "at test_memory_load" + + # Test 3: memory:// load event (hipModuleLoadData). + gdb_test "continue" \ + "Stopped due to GPU code object event.*Inferior loaded memory://\[^\r\n\]+" \ + "load" + + # Test 4: memory:// unload event (hipModuleUnload). + gdb_test "continue" \ + "Stopped due to GPU code object event.*Inferior unloaded memory://.*" \ + "unload" + } + } +} + +with_test_prefix "stop-on-solib-events=0" { + test_disabled +} + +with_test_prefix "stop-on-solib-events=1" { + test_enabled +} -- 2.43.0