From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from simark.ca by simark.ca with LMTP id rw+fDcnvzGqt/zoAWB0awg (envelope-from ) for ; Mon, 12 Oct 2026 10:33:45 -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=AIj/4Wfi; dkim-atps=neutral Received: by simark.ca (Postfix, from userid 112) id 145AE1E051; Mon, 12 Oct 2026 10:33:45 -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 558591E04E for ; Mon, 12 Oct 2026 10:33:41 -0400 (EDT) Received: from vm01.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 26BE34BA23D1 for ; Mon, 12 Oct 2026 14:33:39 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 26BE34BA23D1 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=AIj/4Wfi Received: from SJ2PR03CU001.outbound.protection.outlook.com (mail-westusazon11012003.outbound.protection.outlook.com [52.101.43.3]) by sourceware.org (Postfix) with ESMTPS id 765554BA2E1F for ; Mon, 12 Oct 2026 14:33:02 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 765554BA2E1F 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 765554BA2E1F Authentication-Results: sourceware.org; arc=pass smtp.remote-ip=52.101.43.3 ARC-Seal: i=2; a=rsa-sha256; d=sourceware.org; s=key; t=1791815582; cv=pass; b=LrfenI6Jqyoi0FNeCdXiU40GemmjnqV4P4RvfztpwYyj5Hjc7TAkxCtyHcoysv+jdGB1o3WFyKJ7qWaedEBqOMj7HnnXl3OJaSdWvNLI2mpueiKxmfDsxIC5hbpqq3RFiG6h2200qxQUMA4oroNol8UGQekoA1Gu5Uu1aUB9g3Q= ARC-Message-Signature: i=2; a=rsa-sha256; d=sourceware.org; s=key; t=1791815582; c=relaxed/simple; bh=xyk+kvWm46BhVSYUjUsUVMsGXmqV8SaNVhCKwR+6/Xs=; h=DKIM-Signature:Date:From:To:Subject:Message-ID:MIME-Version; b=KmpGoxTC9g5d+kXPGIvv53j4GyaHAnbMAmDSmCP3qK881DFiLvhym0FUpK4L13rNEQrXcCiLc95gwhiET08QdnEZam8i9DQJpOYTCr670KU4HQmteEEF6ex5bc0oa1dUGf/fBDnSphsR0vHK4lx8l0hU7n3mz7Q94EWymm4B8/s= 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=AIj/4Wfi DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 765554BA2E1F ARC-Seal: i=1; a=rsa-sha256; s=arcselector10001; d=microsoft.com; cv=none; b=qslKyXkoeB/7BFiD0FWSGrOGo+cw8WanPXcyXsuESYN495Wdt2K10Lt62SgV8fR4pFw9LCGHOv8ykioD7uLRyyqgdlfHIVPz1grXQmrcxu2GNvMG5JCpyRUO3fjvyegl50PwskGdZ4b5Mu7IPpbvSw//R1yZq8+v3K+Cl7GihQC8ntDgPuXTE2YnMvhwIn4rqKtTygZMSVeLcR7bg+WWRLPoQqCK78m+yYRxsLKp0ApVGbmG9lKFWEOfNQwS1itNpEpv9BXaF4mQgVkene0m1jLA15jv/dY+ncBGhBgK+1ABaSfqbgmBD5/5CNyTzRrNVpqVdVdv4Q/MUJDnAukxsQ== 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=LiQwGm1NkQB//04XQIkOBPf0iPQ62pMEluczp41nNZQ=; b=g1ByFi6vSC3NbyI0WS2ss0ZxzRXLRAdwy3zZ4cea+OkHNTvjPm3a8RPbqVYdZqHWyBTBtL8s8z5Qg4ZV9mwYlnUjRvJLiOqYUHVgIdSawb5ectL5bKcWq3owoqo+JfgQI/L+uod1sMsqweB3hPsHF0ulWyeKhVTSHCYCp5a9h9aGMousoWvIVqbUnunYQjP2/sb9YkZNZxDdtsd5kbZhJfK1OWtubQno+xmgZyZ9BmaIzJZx6xq4+Zc0s6FWhL82PEp8LmWkAfeZMXrmjXS9OkgnjxL4uM75I50npLs56WHRg+eb5JsuMnWBmEXtYNAbOrU4ZSZ7eN+b8TnE22fHMg== ARC-Authentication-Results: i=1; mx.microsoft.com 1; spf=pass (sender ip is 165.204.84.17) smtp.rcpttodomain=polymtl.ca 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=LiQwGm1NkQB//04XQIkOBPf0iPQ62pMEluczp41nNZQ=; b=AIj/4WfiSIpTYf1G33zfpPgvbFWWdZQZM1Vlc/My2jSw3v4wTFhhkqQiAlhM9hO4tkC67e3k818pujnjGSg4q4I0vyp6Rr2cNRIMIbMVwrfNiKaOLJY1fUWGo5FFf2oUgirPRze9dpbgUbtNBoGC6o+GpQh9x2w9S9HoinSwlQY= Received: from MW4PR03CA0012.namprd03.prod.outlook.com (2603:10b6:303:8f::17) by DSWPR12MB999130.namprd12.prod.outlook.com (2603:10b6:8:36d::16) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_256_GCM_SHA384) id 15.21.496.21; Mon, 12 Oct 2026 14:32:53 +0000 Received: from CO1PEPF000066E8.namprd05.prod.outlook.com (2603:10b6:303:8f:cafe::66) by MW4PR03CA0012.outlook.office365.com (2603:10b6:303:8f::17) with Microsoft SMTP Server (version=TLS1_3, cipher=TLS_AES_256_GCM_SHA384) id 15.21.496.5 via Frontend Transport; Mon, 12 Oct 2026 14:32:53 +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 CO1PEPF000066E8.mail.protection.outlook.com (10.167.249.6) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_256_GCM_SHA384) id 15.21.522.6 via Frontend Transport; Mon, 12 Oct 2026 14:32:53 +0000 Received: from khazad-dum (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.53; Mon, 12 Oct 2026 09:32:46 -0500 Date: Mon, 12 Oct 2026 15:32:36 +0100 From: Lancelot SIX To: Bratislav Filipovic CC: , Subject: Re: [PATCH v3] gdb: Implement stop-on-solib-events for GPU code objects Message-ID: <2j2c6ytzi7ooevyncfwj6siff45zs3loivxc6siuy2tqnjxzrb@logfn6xqblyw> References: <20261008130956.55573-1-bfilipov@amd.com> MIME-Version: 1.0 Content-Type: text/plain; charset="utf-8" Content-Disposition: inline Content-Transfer-Encoding: 8bit In-Reply-To: <20261008130956.55573-1-bfilipov@amd.com> X-Originating-IP: [10.180.168.240] X-ClientProxiedBy: satlexmb08.amd.com (10.181.42.217) To satlexmb07.amd.com (10.181.42.216) X-EOPAttributedMessage: 0 X-MS-PublicTrafficType: Email X-MS-TrafficTypeDiagnostic: CO1PEPF000066E8:EE_|DSWPR12MB999130:EE_ X-MS-Office365-Filtering-Correlation-Id: 03908984-0ef3-4f76-4a2c-08df286db34f X-MS-Exchange-SenderADCheck: 1 X-MS-Exchange-AntiSpam-Relay: 0 X-Microsoft-Antispam: BCL:0; ARA:13230040|1800799024|23010399003|376014|36860700016|82310400026|56012099006|11063799006|5023799004|261009223027099003|10067099003|22082099003|18002099003|6133799003|13003099007|3023799007; X-Microsoft-Antispam-Message-Info: Oy5m7nAUnxEFP08VQuJ7ULQyDC5K+WkhjVFu0PhXZqJ1Kvtkcz837VwoF6IxrbVu9fWEuAV1RBvjlkfd02h93xHx72uh4dRZ71Pc5KIYrVJjpE1i4guB3JEt21WXwNmT2UZJlGvAXVTR/pq4/lvsYLxSW4MrirF+7AbMA3AaYhRI/oKWgRWmj6hhiYInG7nU6ba9SMO9Pag60kOKbv9fcsoAORZCaFT4qR8XG7TITZc4635c3qAmq/4ns6t8PutUf6u1PmgDr7+obnd+3NtxB620xyiCB2ZFb4i1AnmF4FhRPc0QM2L1ilWclJRDOjFtVrNdZXDW/RkG7/0WaCVpYQcRN6dJ8jfCrXDkt1TdSMVzasMbZEKhuCctLAmjYAr6uR9z7HPkH2953U9saLe/6CrIPTXzep1m56GzQ2rVvo5WGb2clXKreS4bG0+Y429DRo1YD+irz8S8lF+Jhb8+1HdYGpZneFdqgDNH6xxslO6E54WluhNRa5gfyKtoT6IM+L1e4qbJpYg58seMM2ZE1vw27O2OZie7vrTzVV/Kz1jCTmWziIeM02yhsSN6vZm6/h+ZCxuSJAk2Z5TOzZUZufggA3GTGeWsRn0G+yqcUlvzbWTvrlHFSLGeC1WbwErXdx6ZK89jfmGnFmxkUwONA7AySAv53+s46WsmYiSCrIg= 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)(1800799024)(23010399003)(376014)(36860700016)(82310400026)(56012099006)(11063799006)(5023799004)(261009223027099003)(10067099003)(22082099003)(18002099003)(6133799003)(13003099007)(3023799007); DIR:OUT; SFP:1101; X-MS-Exchange-AntiSpam-MessageData-ChunkCount: 1 X-MS-Exchange-AntiSpam-MessageData-0: NveNn8fBwoe7p5Ag3FjOmTNVDliJa+zRa3haqYVcRwVmHhDjsz/6PI0pQEFU4iB/8CrRKXCcYzephymWyHuFdNOlbXB/HQfPesXqefITTJ4YrlGYBluz+DvSqssXsCsq3Eh7rjPjl4Pea3C0E2UA83S59H0jk9wH+5m0I8BZLwTy30vGcGkINtIR9vH+izDkO02QJxpql5j5AmAkr9qsXCiFcCcfdh7vUjNPDkV2in7dDVDrFkHx8B7i36jtGIfSJ+ho3i0l1fXxB6tD6X3poMxPMunKPpkJEp6Fywil8TH2ufEVJmfPCGyLHkHBvIAE76eoac1CxN8Q1CvqDYwG48fDM7f91g7HwuyrFJBZaEKyoI7iw5nTwaO27yeqYJevJAMtqk/D9LZTSbG/QzueoA/RLxEnWGDsfGpe2cZC83xUUI4JRkJd88LyCygw4oWr X-OriginatorOrg: amd.com X-MS-Exchange-CrossTenant-OriginalArrivalTime: 12 Oct 2026 14:32:53.3660 (UTC) X-MS-Exchange-CrossTenant-Network-Message-Id: 03908984-0ef3-4f76-4a2c-08df286db34f 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: CO1PEPF000066E8.namprd05.prod.outlook.com X-MS-Exchange-CrossTenant-AuthAs: Anonymous X-MS-Exchange-CrossTenant-FromEntityHeader: HybridOnPrem X-MS-Exchange-Transport-CrossTenantHeadersStamped: DSWPR12MB999130 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 Hi, Thanks for the work you are doing on this. On Thu, Oct 08, 2026 at 03:09:42PM +0200, Bratislav Filipovic wrote: > GDB's "set stop-on-solib-events" 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 does not work 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, 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. Change stop_on_solib_events from 'int' to 'bool' type throughout the > codebase. Update add_setshow_zinteger_cmd to add_setshow_boolean_cmd, > change help text from "If nonzero" to "If enabled", and remove != 0 > comparisons. The command now accepts on/off values while maintaining > backwards compatibility with 0/1. > > 2. Refactor print_solib_event() to use a struct-based design that > centralizes messages and eliminates code duplication between CPU and GPU > event printing. The function now takes a print_solib_event_strings > struct containing translatable message strings and the MI object-kind > field value. This makes it easy to support different types of loadable > objects without repeating GPU code object strings in multiple places. > > 3. Modify check_status() to track when CODE_OBJECT_LIST_UPDATED events > occur during process_event_queue() and update bs->stop and bs->print > based on stop_on_solib_events setting. The bs->print_it field is left > at its default value (print_it_normal), relying on GDB's automatic logic > to set it to print_it_noop when appropriate. This matches the CPU > internal_breakpoint pattern and eliminates redundant assignments. > > 4. Implement print_it() override to display "Stopped due to GPU code object > event" message using the new struct-based print_solib_event design. > > 5. Add MI changes to report object-kind field for all solib events, with > values "shared-library" for CPU libraries and "gpu-code-object" for GPU > code objects. The "library" field name is unchanged for backwards > compatibility. > I am not sure I fully agree we should have different kinds here. At the moment, GPU code objects are reported as any other shared library. When doing `info sharedlibrary`, the GPU code objects are listed the same way as host libraries. Introducing different messages and notifications start to introduce a divergence. I feel like if we wanted to introduce a difference (by solib provider / kind / whatever), this should imply changes beyond print_solib_event to remain consistent. Those are just my thoughts, happy to hear other views. > 6. Add comprehensive test coverage in gdb.rocm/solib-event.exp testing both > file:// events (hipModuleLoad) and memory:// events (hipModuleLoadData), > as well as unload events. The test includes both CLI and MI output > verification. MI tests use manual gdb_expect patterns to avoid triggering > a duplicate notification bug in vanilla GDB's mi_expect_stop. > > When enabled, GDB now stops on GPU code object load/unload events with > "Stopped due to GPU code object event" messages, allowing users to inspect > and debug GPU code object lifecycle consistently with CPU shared libraries. > > --- > Testing notes: > - All 18 tests pass (9 CLI + 9 MI tests) on AMD ROCm systems with gfx90a GPU > - Empirically verified that GPU code object unload events fire automatically > on process termination even without explicit hipModuleUnload calls, > matching CPU shared library cleanup behavior. > - MI tests use manual gdb_expect patterns instead of mi_expect_stop due to > a bug where mi_expect_stop triggers duplicate MI notifications. > - Updated mi-support.exp to accept optional object-kind field in MI > solib-event output for backward compatibility with existing MI consumers. > > gdb/NEWS | 13 + > gdb/amd-dbgapi-target.c | 55 +++- > gdb/break-catch-load.c | 2 +- > gdb/breakpoint.c | 44 +++- > gdb/breakpoint.h | 28 ++- > gdb/doc/gdb.texinfo | 13 +- > gdb/infrun.c | 11 +- > gdb/testsuite/gdb.rocm/solib-event-kernel.cpp | 27 ++ > gdb/testsuite/gdb.rocm/solib-event.cpp | 105 ++++++++ > gdb/testsuite/gdb.rocm/solib-event.exp | 234 ++++++++++++++++++ > gdb/testsuite/lib/mi-support.exp | 2 +- > 11 files changed, 503 insertions(+), 31 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 5d0c4745653..9cd597636d5 100644 > --- a/gdb/NEWS > +++ b/gdb/NEWS > @@ -3,6 +3,11 @@ > > *** 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, 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 > @@ -79,6 +84,14 @@ info locals > The new "-shadowed on|off" option overrides the "set print shadowed" > setting for a single invocation of the command. > > +* MI changes > + > +** The solib-event notification now includes an 'object-kind' field > + to distinguish different types of loadable objects. Currently > + defined values are 'shared-library' for CPU shared libraries and > + 'gpu-code-object' for GPU code objects. For backwards compatibility, > + the 'library' field name is used for both types of objects. > + > *** Changes in GDB 18 > > * Support for the Common Trace Format (CTF) has been removed. GDB now > diff --git a/gdb/amd-dbgapi-target.c b/gdb/amd-dbgapi-target.c > index 93207a8fdd2..64e9a3aee94 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 > > @@ -275,7 +277,17 @@ struct amd_dbgapi_inferior_info > wave_info_map; > }; > > -static amd_dbgapi_event_id_t process_event_queue > +/* Return value for process_event_queue. */ > +struct process_event_queue_result > +{ > + /* The event ID, or AMD_DBGAPI_EVENT_NONE if no event was found. */ > + amd_dbgapi_event_id_t event_id; > + > + /* True if a CODE_OBJECT_LIST_UPDATED event was seen during processing. */ > + bool code_object_list_updated; > +}; > + > +static process_event_queue_result process_event_queue > (amd_dbgapi_inferior_info &info, > amd_dbgapi_event_kind_t until_event_kind = AMD_DBGAPI_EVENT_KIND_NONE); > > @@ -587,6 +599,7 @@ struct amd_dbgapi_target_breakpoint : public code_breakpoint > > void re_set (program_space *) override; > void check_status (struct bpstat *bs) override; > + enum print_stop_action print_it (const bpstat *bs) const override; > }; > > void > @@ -603,7 +616,6 @@ amd_dbgapi_target_breakpoint::check_status (struct bpstat *bs) > amd_dbgapi_status_t status; > > bs->stop = 0; > - bs->print_it = print_it_noop; > > /* Find the address the breakpoint is set at. */ > auto match_breakpoint > @@ -639,7 +651,7 @@ amd_dbgapi_target_breakpoint::check_status (struct bpstat *bs) > > /* If the action is AMD_DBGAPI_BREAKPOINT_ACTION_HALT, we need to wait until > a breakpoint resume event for this breakpoint_id is seen. */ > - amd_dbgapi_event_id_t resume_event_id > + auto [resume_event_id, code_object_list_updated] > = process_event_queue (info, AMD_DBGAPI_EVENT_KIND_BREAKPOINT_RESUME); > > /* We should always get a breakpoint_resume event after processing all > @@ -665,6 +677,26 @@ amd_dbgapi_target_breakpoint::check_status (struct bpstat *bs) > pulongest (resume_breakpoint_id.handle)); > > amd_dbgapi_event_processed (resume_event_id); > + > + /* If a CODE_OBJECT_LIST_UPDATED event was seen during event processing, > + check if the user requested to stop on solib events. This implements > + stop-on-solib-events for GPU code objects. */ > + if (code_object_list_updated && stop_on_solib_events) > + { > + bs->stop = true; > + bs->print = true; > + } > +} > + > +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. */ > + print_solib_event (false, print_solib_event_strings_gpu_code_object); > + > + return PRINT_NOTHING; > } > > bool > @@ -1460,7 +1492,8 @@ add_gpu_thread (inferior *inf, ptid_t wave_ptid) > static void > process_one_event (amd_dbgapi_inferior_info &info, > amd_dbgapi_event_id_t event_id, > - amd_dbgapi_event_kind_t event_kind) > + amd_dbgapi_event_kind_t event_kind, > + bool &code_object_list_updated) > { > /* Automatically mark this event processed when going out of scope. */ > scoped_amd_dbgapi_event_processed mark_event_processed (event_id); > @@ -1565,6 +1598,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 ()); > + code_object_list_updated = true; > handle_solib_event (); > break; > > @@ -1639,7 +1673,7 @@ event_kind_str (amd_dbgapi_event_kind_t kind) > Wave stop events that are not returned are queued into their inferior's > amd_dbgapi_inferior_info pending wave events. */ > > -static amd_dbgapi_event_id_t > +static process_event_queue_result > process_event_queue (amd_dbgapi_inferior_info &info, > amd_dbgapi_event_kind_t until_event_kind) > { > @@ -1647,6 +1681,8 @@ process_event_queue (amd_dbgapi_inferior_info &info, > performance, make sure it is not required. */ > gdb_assert (!info.forward_progress_required); > > + bool code_object_list_updated = false; > + > while (true) > { > amd_dbgapi_event_id_t event_id; > @@ -1667,9 +1703,9 @@ process_event_queue (amd_dbgapi_inferior_info &info, > event_kind_str (event_kind)); > > if (event_id == AMD_DBGAPI_EVENT_NONE || event_kind == until_event_kind) > - return event_id; > + return {event_id, code_object_list_updated}; > > - process_one_event (info, event_id, event_kind); > + process_one_event (info, event_id, event_kind, code_object_list_updated); > } > } > > @@ -1809,6 +1845,11 @@ amd_dbgapi_target::wait (ptid_t ptid, struct target_waitstatus *ws, > } > } > > + /* We are returning a real stop event, so don't mark the async handler. > + This prevents the event loop from re-processing the same stop event, > + which would cause duplicate MI notifications. */ > + more_events.release (); > + This bit seems odd. We purposefully mark the event handler so amd_dbgapi_target::wait gets called until it returns minus_one_ptid. Doing this release could mean that we still have events pending (either in wave_events, or down in dbgapi / the driver), but GDB's core might not know and wait to be notified, which might not happen. Best, Lancelot. > *ws = gpu_waitstatus; > return event_ptid; > } > diff --git a/gdb/break-catch-load.c b/gdb/break-catch-load.c > index 5492a56b423..b4e0a2978cb 100644 > --- a/gdb/break-catch-load.c > +++ b/gdb/break-catch-load.c > @@ -151,7 +151,7 @@ solib_catchpoint::print_it (const bpstat *bs) const > uiout->text ("\n"); > if (uiout->is_mi_like_p ()) > uiout->field_string ("disp", bpdisp_text (this->disposition)); > - print_solib_event (true); > + print_solib_event (true, print_solib_event_strings_shared_library); > return PRINT_SRC_AND_LOC; > } > > diff --git a/gdb/breakpoint.c b/gdb/breakpoint.c > index 0f5360c4ce5..e2c6614b68d 100644 > --- a/gdb/breakpoint.c > +++ b/gdb/breakpoint.c > @@ -5368,24 +5368,44 @@ print_bp_stop_message (bpstat *bs) > > /* See breakpoint.h. */ > > +const print_solib_event_strings print_solib_event_strings_shared_library = > +{ > + N_("Stopped due to shared library event:\n"), > + N_("Stopped due to shared library event (no libraries added or " > + "removed)\n"), > + "shared-library" > +}; > + > +/* See breakpoint.h. */ > + > +const print_solib_event_strings print_solib_event_strings_gpu_code_object = > +{ > + N_("Stopped due to GPU code object event:\n"), > + N_("Stopped due to GPU code object event (no GPU code objects added or " > + "removed)\n"), > + "gpu-code-object" > +}; > + > +/* See breakpoint.h. */ > + > void > -print_solib_event (bool is_catchpoint) > +print_solib_event (bool is_catchpoint, > + const print_solib_event_strings &strings) > { > bool any_deleted = !current_program_space->deleted_solibs.empty (); > bool any_added = !current_program_space->added_solibs.empty (); > > if (!is_catchpoint) > - { > - if (any_added || any_deleted) > - current_uiout->text (_("Stopped due to shared library event:\n")); > - else > - current_uiout->text (_("Stopped due to shared library event (no " > - "libraries added or removed)\n")); > - } > + current_uiout->text (any_added || any_deleted > + ? _(strings.stopped_msg) > + : _(strings.stopped_no_change_msg)); > > 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)); > + current_uiout->field_string ("object-kind", strings.object_kind); > + } > > if (any_deleted) > { > @@ -5462,7 +5482,7 @@ bpstat_print (bpstat *bs, target_waitkind kind) > OS-level shared library event, do the same thing. */ > if (kind == TARGET_WAITKIND_LOADED) > { > - print_solib_event (false); > + print_solib_event (false, print_solib_event_strings_shared_library); > return PRINT_NOTHING; > } > > @@ -12547,7 +12567,7 @@ internal_breakpoint::print_it (const bpstat *bs) const > /* Did we stop because the user set the stop_on_solib_events > variable? (If so, we report this as a generic, "Stopped due > to shlib event" message.) */ > - print_solib_event (false); > + print_solib_event (false, print_solib_event_strings_shared_library); > break; > > case bp_thread_event: > diff --git a/gdb/breakpoint.h b/gdb/breakpoint.h > index ba6eb15a47e..5cc5d814a54 100644 > --- a/gdb/breakpoint.h > +++ b/gdb/breakpoint.h > @@ -2098,11 +2098,37 @@ extern void catch_exception_event (enum exception_event_kind ex_event, > const char *regex, bool tempflag, > int from_tty); > > +/* Strings used by print_solib_event to describe a kind of loadable > + object. */ > + > +struct print_solib_event_strings > +{ > + /* Message printed when objects were added or removed. */ > + const char *stopped_msg; > + > + /* Message printed when no objects were added or removed. */ > + const char *stopped_no_change_msg; > + > + /* Value of the MI "object-kind" field. */ > + const char *object_kind; > +}; > + > +/* Strings for shared libraries. */ > + > +extern const print_solib_event_strings > + print_solib_event_strings_shared_library; > + > +/* Strings for GPU code objects. */ > + > +extern const print_solib_event_strings > + print_solib_event_strings_gpu_code_object; > + > /* A helper function that prints a shared library stopped 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 print_solib_event_strings &strings); > > /* 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 4bd5357c587..397f0916a6d 100644 > --- a/gdb/doc/gdb.texinfo > +++ b/gdb/doc/gdb.texinfo > @@ -22536,14 +22536,15 @@ The surrounding square brackets are optional. > @item set stop-on-solib-events @r{[}on@r{|}off@r{]} > @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. > +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 > @@ -32695,6 +32696,10 @@ The inferior has stopped due to a library being loaded or unloaded. > This can happen when @code{stop-on-solib-events} (@pxref{Files}) is > set or when a @code{catch load} or @code{catch unload} catchpoint is > in use (@pxref{Set Catchpoints}). > + > +This stop reason may include an optional @code{object-kind} field > +to distinguish different types of loadable objects. Currently defined > +values are @code{gpu-code-object} for GPU code objects. > @item fork > The inferior has forked. This is reported when @code{catch fork} > (@pxref{Set Catchpoints}) has been used. > diff --git a/gdb/infrun.c b/gdb/infrun.c > index 47966cf72db..3492821f4a9 100644 > --- a/gdb/infrun.c > +++ b/gdb/infrun.c > @@ -413,8 +413,8 @@ get_detach_signal (process_stratum_target *proc_target, ptid_t ptid) > > static struct cmd_list_element *stop_command; > > -/* See infrun.h. */ > - > +/* True if we want to give control to the user when we're notified > + of shared library events by the dynamic linker. */ > bool stop_on_solib_events; > > /* Enable or disable optional shared library event breakpoints > @@ -10850,9 +10850,10 @@ leave it stopped or free to run as needed."), > &stop_on_solib_events, _("\ > Set stopping for shared library events."), _("\ > Show stopping for shared library events."), _("\ > -If on, 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."), > +If enabled, gdb will give control to the user when the dynamic linker\n\ > +notifies gdb of shared library events, or when GPU code objects are loaded\n\ > +or unloaded. The most common events of interest to the user would be\n\ > +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..87f12ccdb46 > --- /dev/null > +++ b/gdb/testsuite/gdb.rocm/solib-event.exp > @@ -0,0 +1,234 @@ > +# 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 . > + > +# Test that "set stop-on-solib-events on" 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 (off) does not stop execution. > +# Also tests MI output includes object-kind field. > + > +load_lib rocm.exp > +load_lib mi-support.exp > + > +require allow_hip_tests > + > +standard_testfile .cpp > + > +# Build GPU code object. > +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 on" > + > + 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://\[^\r\n\]+" \ > + "unload" > + } > + } > +} > + > +proc test_mi_output {} { > + global MIFLAGS > + global hipmodule_path > + global mi_gdb_prompt > + > + set MIFLAGS "-i=mi2" > + > + with_rocm_gpu_lock { > + if {[mi_clean_restart $::testfile]} { > + return > + } > + > + # Set inferior arguments before running to main. > + mi_gdb_test "-gdb-set args $hipmodule_path" "\\^done" \ > + "set inferior arguments" > + > + if {![mi_runto_main]} { > + return > + } > + > + mi_create_breakpoint [gdb_get_line_number "Enable SOLIB events here"] \ > + "insert breakpoint at enable marker" > + > + mi_send_resuming_command "exec-continue" "continue to enable marker" > + # The warm_up<<<>>>() kernel loads code objects here, emitting > + # async =library-loaded notifications even though stop-on-solib-events=off. > + # Consume these async notifications before expecting the breakpoint stop. > + gdb_expect { > + -re "=library-loaded\[^\r\n\]*\r\n" { > + exp_continue > + } > + -re "\\*stopped,reason=\"breakpoint-hit\"\[^\r\n\]*\r\n$mi_gdb_prompt" { > + pass "stop at enable marker" > + } > + timeout { > + fail "stop at enable marker (timeout)" > + } > + } > + > + mi_gdb_test "111-gdb-set stop-on-solib-events on" "111\\^done" \ > + "set stop-on-solib-events via MI" > + > + mi_create_breakpoint "test_file_load" "insert breakpoint at test_file_load" > + > + # Use manual gdb_expect instead of mi_expect_stop to avoid triggering > + # duplicate MI notification bug in vanilla GDB. > + mi_send_resuming_command "exec-continue" "continue to test_file_load" > + > + set breakpoint_modified_count 0 > + set console_output_count 0 > + set stopped_count 0 > + > + gdb_expect { > + -re "=breakpoint-modified,bkpt=\{number=\"\[0-9\]+\"" { > + incr breakpoint_modified_count > + exp_continue > + } > + -re "~\"\\\\n\"\r\n" { > + incr console_output_count > + exp_continue > + } > + -re "~\"Thread \[0-9\]+ .*hit Breakpoint \[0-9\]+, test_file_load" { > + exp_continue > + } > + -re "~\[^\r\n\]*\r\n" { > + exp_continue > + } > + -re "\\*stopped,reason=\"breakpoint-hit\".*func=\"test_file_load\"\[^\r\n\]*\r\n$mi_gdb_prompt" { > + incr stopped_count > + if {$breakpoint_modified_count == 1 && $stopped_count == 1} { > + pass "stop at test_file_load" > + } else { > + fail "stop at test_file_load (breakpoint-modified=$breakpoint_modified_count, stopped=$stopped_count)" > + } > + } > + timeout { > + fail "stop at test_file_load (timeout)" > + } > + } > + > + # Check for GPU code object load event with object-kind field. > + mi_send_resuming_command "exec-continue" "continue for GPU load event" > + > + set library_loaded_count 0 > + set solib_stopped_count 0 > + set has_object_kind 0 > + > + gdb_expect { > + -re "=library-loaded\[^\r\n\]*\r\n" { > + incr library_loaded_count > + exp_continue > + } > + -re "~\"Stopped due to GPU code object event" { > + exp_continue > + } > + -re "~\[^\r\n\]*\r\n" { > + exp_continue > + } > + -re "\\*stopped,reason=\"solib-event\".*object-kind=\"gpu-code-object\".*library=\[^\r\n\]*\r\n" { > + incr solib_stopped_count > + set has_object_kind 1 > + exp_continue > + } > + -re "$mi_gdb_prompt" { > + if {$solib_stopped_count == 1 && $has_object_kind == 1} { > + pass "verify object-kind field in MI output" > + } else { > + fail "verify object-kind field in MI output (stopped=$solib_stopped_count, has_object_kind=$has_object_kind)" > + } > + } > + timeout { > + fail "verify object-kind field in MI output (timeout)" > + } > + } > + } > +} > + > +with_test_prefix "stop-on-solib-events=off" { > + test_disabled > +} > + > +with_test_prefix "stop-on-solib-events=on" { > + test_enabled > +} > + > +with_test_prefix "MI output" { > + test_mi_output > +} > diff --git a/gdb/testsuite/lib/mi-support.exp b/gdb/testsuite/lib/mi-support.exp > index cd7e0a61189..20f2a9bbc96 100644 > --- a/gdb/testsuite/lib/mi-support.exp > +++ b/gdb/testsuite/lib/mi-support.exp > @@ -1250,7 +1250,7 @@ proc mi_expect_stop { reason func args file line extra test } { > } > > if { $reason == "solib-event" } { > - set pattern "\\*stopped,reason=\"solib-event\",thread-id=\"$decimal\",stopped-threads=$any\r\n($thread_selected_re|$breakpoint_re)*$prompt_re" > + set pattern "\\*stopped,reason=\"solib-event\"(,object-kind=\"\[^\"]*\")?,thread-id=\"$decimal\",stopped-threads=$any\r\n($thread_selected_re|$breakpoint_re)*$prompt_re" > verbose -log "mi_expect_stop: expecting: $pattern" > gdb_expect { > -re "$pattern" { > > base-commit: 60b08e6b3adf81fbd95cc69511cac640714174e2 > -- > 2.43.0