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

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