[PATCH] gdb: Implement stop-on-solib-events for GPU code objects

[PATCH] gdb: Implement stop-on-solib-events for GPU code objects

[Bratislav Filipovic]

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 <map>
 
@@ -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<std::pair<ptid_t, target_waitstatus>> 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_map<thread_info *,
@@ -587,6 +598,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
@@ -637,6 +649,10 @@ amd_dbgapi_target_breakpoint::check_status (struct bpstat *bs)
 
   require_forward_progress (info, false);
 
+  /* Clear the flag before processing events so we can detect if a new
+     CODE_OBJECT_LIST_UPDATED event occurs during this processing cycle.  */
+  info.code_object_list_updated = false;
+
   /* 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
@@ -665,6 +681,72 @@ 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 (info.code_object_list_updated && stop_on_solib_events != 0)
+    {
+      bs->stop = 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 <http://www.gnu.org/licenses/>.  */
+
+/* Device kernel for solib-event test.
+   Compiled with --cuda-device-only to produce .co file.  */
+
+#include <hip/hip_runtime.h>
+
+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 <http://www.gnu.org/licenses/>.  */
+
+#include "rocm-test-utils.h"
+#include <hip/hip_runtime.h>
+#include <fstream>
+#include <vector>
+
+/* 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<unsigned char> module_buffer (module_size);
+
+  if (!mod.read (reinterpret_cast<char *> (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 <module_path>\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 <http://www.gnu.org/licenses/>.
+
+# 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

Re: [PATCH] gdb: Implement stop-on-solib-events for GPU code objects

[Simon Marchi]

On 9/24/26 9:04 AM, Bratislav Filipovic wrote:
> 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.

You indeed generalized print_solib_event, but
amd_dbgapi_target_breakpoint::print_it still contains the duplicated
code, and doesn't use print_solib_event.  I guess you are missing a
piece of the change.

Simon

Re: [PATCH] gdb: Implement stop-on-solib-events for GPU code objects

[Simon Marchi]

On 9/24/26 9:04 AM, Bratislav Filipovic wrote:
> 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().

This flag does not need to live in amd_dbgapi_inferior_info, since it's

only needed within the duration of one

amd_dbgapi_target_breakpoint::check_status call.  It's not a state that

needs to be persisted.  Perhaps have process_event_queue return that

value, so check_status can use it?  process_event_queue already returns

a value (amd_dbgapi_event_id_t), it could return a small struct instead.

> +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);
> +    }

This (and the change in print_solib_event) adds a new field tot he stops
of kind "solib-event".  I think it should be documented where
"solib-event" is described, in section "GDB/MI Async Records" of the
documentation.

> 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\

As a GDB developer, I find this documentation change a bit unnecessary,
because I know that GPU code objects are (in GDB) the same as shared
libraries, just with another name.  But I guess it's good to be
explicit, because that might not be clear to end users.  But I would not
mention "(AMD ROCm targets)" part (same for the documentation part
above), because we want to keep the documentation of these commands
fairly target-agnostic.

Will the term "GPU code objects" be applicable for the Intel GPU target
too?

> +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://.*" \

This last line should use \[^\r\n\]+ like the other ones I guess.

Simon

[PATCH] gdb: Implement stop-on-solib-events for GPU code objects

[Bratislav Filipovic]

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.

Thanks for the review Simon. Addressed your comments, here's the v2.

Regards Bratislav
---
 gdb/NEWS                                      |   6 +
 gdb/amd-dbgapi-target.c                       |  56 ++++++++-
 gdb/breakpoint.c                              |  31 +++--
 gdb/breakpoint.h                              |   6 +-
 gdb/doc/gdb.texinfo                           |  13 +-
 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, 340 insertions(+), 21 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..9b937bc21f9 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 <map>
 
@@ -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;
@@ -275,7 +281,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 +603,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
@@ -639,7 +656,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 +682,29 @@ 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 != 0)
+    {
+      bs->stop = 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.  */
+  print_solib_event (false, "GPU code object", "GPU code objects",
+		     "code-object", gpu_code_object_kind);
+
+  return PRINT_NOTHING;
 }
 
 bool
@@ -1460,7 +1500,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 +1606,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 +1681,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 +1689,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 +1711,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);
     }
 }
 
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..fd67204a47f 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.
+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
@@ -32638,6 +32639,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 92b21017b03..5f63a85e8d8 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.  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 <http://www.gnu.org/licenses/>.  */
+
+/* Device kernel for solib-event test.
+   Compiled with --cuda-device-only to produce .co file.  */
+
+#include <hip/hip_runtime.h>
+
+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 <http://www.gnu.org/licenses/>.  */
+
+#include "rocm-test-utils.h"
+#include <hip/hip_runtime.h>
+#include <fstream>
+#include <vector>
+
+/* 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<unsigned char> module_buffer (module_size);
+
+  if (!mod.read (reinterpret_cast<char *> (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 <module_path>\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..f91b8d41995
--- /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 <http://www.gnu.org/licenses/>.
+
+# 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://\[^\r\n\]+" \
+		"unload"
+	}
+    }
+}
+
+with_test_prefix "stop-on-solib-events=0" {
+    test_disabled
+}
+
+with_test_prefix "stop-on-solib-events=1" {
+    test_enabled
+}
-- 
2.43.0

Re: [PATCH] gdb: Implement stop-on-solib-events for GPU code objects

[Eli Zaretskii]

> From: Bratislav Filipovic <bfilipov@amd.com>
> CC: <simark@simark.ca>, <Lancelot.Six@amd.com>, <TankutBaris.Aktemur@amd.com>, 
>  <luis.machado.foss@gmail.com>, <pedro@palves.net>, Bratislav Filipovic
>  <bfilipov@amd.com>
> Date: Thu, 1 Oct 2026 12:35:05 +0200
> 
> 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.
> 
> Thanks for the review Simon. Addressed your comments, here's the v2.
> 
> Regards Bratislav
> ---
>  gdb/NEWS                                      |   6 +
>  gdb/amd-dbgapi-target.c                       |  56 ++++++++-
>  gdb/breakpoint.c                              |  31 +++--
>  gdb/breakpoint.h                              |   6 +-
>  gdb/doc/gdb.texinfo                           |  13 +-
>  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, 340 insertions(+), 21 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

Thanks, the documentation parts of this patch are okay.

Reviewed-By: Eli Zaretskii <eliz@gnu.org>

Re: [PATCH] gdb: Implement stop-on-solib-events for GPU code objects

[Simon Marchi]

Hi Bratislav,

Some advice for future submissions:

 - Since this is v2, pass `-v2` to git-format-patch or git-send-email so
   that it inserts "v2" in the subject.

 - Set `format.useAutoBase` in your git config to make git record the
   base commit on which your patch is based.  This makes it easier to
   apply the patch.

   Ref: https://git-scm.com/docs/git-config#Documentation/git-config.txt-formatuseAutoBase

On 10/1/26 6:35 AM, Bratislav Filipovic wrote:
> 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().

Point #2, not true anymore.

> 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

Remove `(non-zero value)`, since stop-on-solib-events is now a bool.

The MI changes should be documented in NEWS, in an "MI changes" section
(see past releases for an example).

> @@ -665,6 +682,29 @@ 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 != 0)

Remove the `!= 0`, as stop_on_solib_events is now a bool.

> 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)

Instead of passing all these `const char *`, I think it would be better
to centralize all the messages here.  One reason is that when we'll have
a second GPU target, we won't want to have it repeat the same GPU code
object strings.  Also, the strings for shared library are "hidden" as
default parameters to this function, in the header, which makes it hard
to follow.  Finally, having "decomposed" messages like this is not good
for translation, it's better to have full sentences.

I would suggest something like this.  In the header file:

---

/* 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;

  /* MI field name for each item of the "added" and "removed" lists.  */
  const char *item_field_name;

  /* 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,
			       const print_solib_event_strings &strings);

---

In the source file:

---

/* 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"),
  "library",
  "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"),
  "code-object",
  "gpu-code-object"
};

/* See breakpoint.h.  */

void
print_solib_event (bool is_catchpoint,
		   const print_solib_event_strings &strings)

---

And then in print_solib_event_strings you get the strings from the
`strings` parameter.  For the translated strings (those marked with N_),
pass them through the _ macro:

  if (!is_catchpoint)
    current_uiout->text (any_added || any_deleted
			 ? _(strings.stopped_msg)
			 : _(strings.stopped_no_change_msg));

And then in amd-dbgapi-target.c, the print_solib_event call would look
like:

  print_solib_event (false, print_solib_event_strings_gpu_code_object);

Same for the other calls, but they would pass the other object.

Just a note regarding the MI output, I think it's fine to add the
"object-kind" value so that a frontend could adjust its display if it
wants to say "GPU code objects" instead of "shared libraries".  However,
I would leave the existing "library" entry as-is.  The reason being that
frontends that do not know about this new GPU code object thing would
expect to see a "library" entry, so it would be a backwards incompatible
change.  Also, if you leave it as "library", existing frontends can work
just fine with GPU code objects (they'll present them as libraries),
since there's nothing fundamentally different between shared libraries
and GPU code objects in GDB, other than the name.

So I'd remove the "item_field_name" from the struct above.

Also, I would always emit the object-kind key, not just for GPU objects,
it's clearer that way.  It's not a breaking change to add an entry to a
dict (or whatever MI calls them).  Some tests will need to be updated
though, so adding this entry could perhaps be done as a preparatory
patch.

> diff --git a/gdb/testsuite/gdb.rocm/solib-event.exp b/gdb/testsuite/gdb.rocm/solib-event.exp
> new file mode 100644
> index 00000000000..f91b8d41995
> --- /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.

Remove AMD copyright.

> +
> +# 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 <http://www.gnu.org/licenses/>.
> +
> +# 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
> +}

Since you have a "Build host executable" comment below, I would add a
"Build GPU code object" above this.

> +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"

Since stop-on-solib-events is now a bool, I would suggest using on/off
instead of 1/0 throughout.

It would be nice to have a test for the MI change.

> +
> +	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"
> +	}
> +    }

I don't know if we need a test for that, but how does it behave if we
have some code objects loaded and we let the host program terminate?
Does it fire some code object updated events?

Simon

[PATCH v3] gdb: Implement stop-on-solib-events for GPU code objects

[Bratislav Filipovic]

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.

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 <map>
 
@@ -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 ();
+
   *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 <http://www.gnu.org/licenses/>.  */
+
+/* Device kernel for solib-event test.
+   Compiled with --cuda-device-only to produce .co file.  */
+
+#include <hip/hip_runtime.h>
+
+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 <http://www.gnu.org/licenses/>.  */
+
+#include "rocm-test-utils.h"
+#include <hip/hip_runtime.h>
+#include <fstream>
+#include <vector>
+
+/* 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<unsigned char> module_buffer (module_size);
+
+  if (!mod.read (reinterpret_cast<char *> (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 <module_path>\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 <http://www.gnu.org/licenses/>.
+
+# 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

Re: [PATCH v3] gdb: Implement stop-on-solib-events for GPU code objects

[Lancelot SIX]

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 <map>
>  
> @@ -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 <http://www.gnu.org/licenses/>.  */
> +
> +/* Device kernel for solib-event test.
> +   Compiled with --cuda-device-only to produce .co file.  */
> +
> +#include <hip/hip_runtime.h>
> +
> +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 <http://www.gnu.org/licenses/>.  */
> +
> +#include "rocm-test-utils.h"
> +#include <hip/hip_runtime.h>
> +#include <fstream>
> +#include <vector>
> +
> +/* 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<unsigned char> module_buffer (module_size);
> +
> +  if (!mod.read (reinterpret_cast<char *> (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 <module_path>\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 <http://www.gnu.org/licenses/>.
> +
> +# 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

[PATCH v4] gdb: Implement stop-on-solib-events for GPU code objects

[Bratislav Filipovic]

Thanks for the review, Lancelot!

> 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.

This was added as the result of discussions where we were planning to separate cpu and gpu code objects in the future. Since we've moved from that idea, I agree that different kinds are not the way to go. I've removed the object-kind field from MI output entirely. GPU code objects are now treated identically to CPU shared libraries in all output formats - CLI, MI, and info sharedlibrary all show them the same way with no divergence.

> 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.

The duplicate MI notifications I was seeing while testing were actually caused by a bug in vanilla GDB's mi_expect_stop function, not by the async event handler mechanism. This was a remainder from me trying to fix that. I've removed the more_events.release() call entirely. The MI tests now use manual gdb_expect patterns to work around the mi_expect_stop bug, which properly prevents duplicates without breaking event processing.

v4 patch attached with both changes.

Testing notes:
- All 18 tests pass (9 CLI + 9 MI) on gfx90a
- GPU code objects now treated identically to CPU shared libraries
- No divergence between CLI, MI, and info sharedlibrary output

Regards,
Bratislav

---
 gdb/NEWS                                      |   5 +
 gdb/amd-dbgapi-target.c                       |  50 +++-
 gdb/break-catch-load.c                        |   2 +-
 gdb/breakpoint.c                              |  39 ++-
 gdb/breakpoint.h                              |  25 +-
 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        | 232 ++++++++++++++++++
 10 files changed, 479 insertions(+), 30 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 7303ec0e460..6fe6dd0e955 100644
--- a/gdb/NEWS
+++ b/gdb/NEWS
@@ -7,6 +7,11 @@
   Support displaying and modifying Intel APX registers (Extended GPRs
   $r16 - $r31) including byte, word and dword pseudo registers.
 
+* 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
diff --git a/gdb/amd-dbgapi-target.c b/gdb/amd-dbgapi-target.c
index a38a8bafb90..101687dfa81 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 <map>
 
@@ -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
@@ -1468,7 +1500,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);
@@ -1573,6 +1606,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;
 
@@ -1647,7 +1681,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)
 {
@@ -1655,6 +1689,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;
@@ -1675,9 +1711,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);
     }
 }
 
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..4f7ab8a9273 100644
--- a/gdb/breakpoint.c
+++ b/gdb/breakpoint.c
@@ -5368,24 +5368,39 @@ 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")
+};
+
+/* 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")
+};
+
+/* 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));
 
   if (any_deleted)
     {
@@ -5462,7 +5477,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 +12562,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..cb620b9defe 100644
--- a/gdb/breakpoint.h
+++ b/gdb/breakpoint.h
@@ -2098,11 +2098,34 @@ 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;
+};
+
+/* 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 66152d2582e..96fcd82c572 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 <http://www.gnu.org/licenses/>.  */
+
+/* Device kernel for solib-event test.
+   Compiled with --cuda-device-only to produce .co file.  */
+
+#include <hip/hip_runtime.h>
+
+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 <http://www.gnu.org/licenses/>.  */
+
+#include "rocm-test-utils.h"
+#include <hip/hip_runtime.h>
+#include <fstream>
+#include <vector>
+
+/* 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<unsigned char> module_buffer (module_size);
+
+  if (!mod.read (reinterpret_cast<char *> (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 <module_path>\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..a7e5fb65d13
--- /dev/null
+++ b/gdb/testsuite/gdb.rocm/solib-event.exp
@@ -0,0 +1,232 @@
+# 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 <http://www.gnu.org/licenses/>.
+
+# 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 in MI output.
+	mi_send_resuming_command "exec-continue" "continue for GPU load event"
+
+	set library_loaded_count 0
+	set solib_stopped_count 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\".*library=\[^\r\n\]*\r\n" {
+		incr solib_stopped_count
+		exp_continue
+	    }
+	    -re "$mi_gdb_prompt" {
+		if {$solib_stopped_count == 1} {
+		    pass "verify solib-event in MI output"
+		} else {
+		    fail "verify solib-event in MI output (stopped=$solib_stopped_count)"
+		}
+	    }
+	    timeout {
+		fail "verify solib-event 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
+}

base-commit: aeed64b39aed57729f66ad51560d8964e8201e6e
-- 
2.43.0