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

gdb/amd-dbgapi-target.c

@@ -38,13 +38,15 @@
 #include "gdbthread.h"
 #include "inf-loop.h"
 #include "inferior.h"
+#include "infrun.h"
 #include "location.h"
 #include "objfiles.h"
 #include "observable.h"
 #include "registry.h"
 #include "solib.h"
 #include "target.h"
 #include "tid-parse.h"
+#include "ui-out.h"
 #include "bfd/elf-bfd.h"
 #include "elf/amdgpu.h"
 #include "cp-support.h"
@@ -284,6 +286,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 *,
@@ -832,6 +839,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
@@ -882,6 +890,12 @@ amd_dbgapi_target_breakpoint::check_status (struct bpstat *bs)
 
   require_forward_progress (info, false);
 
+  /* Clear the flag so event processing below can detect a new
+     CODE_OBJECT_LIST_UPDATED event.  This flag is reset at the top of
+     check_status before each use to ensure we only see events from the
+     current stop.  */
+  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
@@ -910,6 +924,74 @@ 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 GPU code object events.  This
+     implements stop-on-solib-events for GPU code objects.  */
+  if (info.code_object_list_updated
+      && (stop_on_solib_events == STOP_SOLIB_GPU
+	  || stop_on_solib_events == STOP_SOLIB_ALL))
+    {
+      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 set to STOP_SOLIB_GPU or STOP_SOLIB_ALL.  */
+  bool any_deleted = !current_program_space->deleted_solibs.empty ();
+  bool any_added = !current_program_space->added_solibs.empty ();
+
+  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");
+    }
+
+  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 (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
@@ -2120,6 +2202,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;
 

gdb/breakpoint.c

@@ -12392,8 +12392,10 @@ internal_breakpoint::check_status (bpstat *bs)
 	 events.  This allows the user to get control and place
 	 breakpoints in initializer routines for dynamically loaded
 	 objects (among other things).  */
-      bs->stop = stop_on_solib_events != 0;
-      bs->print = stop_on_solib_events != 0;
+      bs->stop = (stop_on_solib_events == STOP_SOLIB_CPU
+		  || stop_on_solib_events == STOP_SOLIB_ALL);
+      bs->print = (stop_on_solib_events == STOP_SOLIB_CPU
+		   || stop_on_solib_events == STOP_SOLIB_ALL);
     }
   else
     bs->stop = false;

gdb/infrun.c

@@ -434,6 +434,17 @@ static struct cmd_list_element *stop_command;
    of shared library events by the dynamic linker.  */
 int stop_on_solib_events;
 
+/* Extra literals for stop-on-solib-events setting.
+   Allows both string keywords and numeric values (0-3).  */
+static const literal_def stop_on_solib_events_literals[] =
+  {
+    { "none", STOP_SOLIB_NONE, 0 },  /* Accept "none" or 0.  */
+    { "cpu",  STOP_SOLIB_CPU,  1 },  /* Accept "cpu" or 1.  */
+    { "gpu",  STOP_SOLIB_GPU,  2 },  /* Accept "gpu" or 2.  */
+    { "all",  STOP_SOLIB_ALL,  3 },  /* Accept "all" or 3.  */
+    { nullptr }
+  };
+
 /* Enable or disable optional shared library event breakpoints
    as appropriate when the above flag is changed.  */
 
@@ -448,8 +459,41 @@ static void
 show_stop_on_solib_events (struct ui_file *file, int from_tty,
 			   struct cmd_list_element *c, const char *value)
 {
-  gdb_printf (file, _("Stopping for shared library events is %s.\n"),
-	      value);
+  /* The value string will show the literal ("none", "cpu", "gpu", "all")
+     when set via keyword, or the number when set numerically.  Access the
+     actual integer value to provide consistent descriptive suffix.  */
+  switch (stop_on_solib_events)
+    {
+    case STOP_SOLIB_NONE:
+      gdb_printf (file,
+		  _("Stopping for shared library events is \"%s\" "
+		    "(disabled).\n"),
+		  value);
+      break;
+    case STOP_SOLIB_CPU:
+      gdb_printf (file,
+		  _("Stopping for shared library events is \"%s\" "
+		    "(CPU libraries only).\n"),
+		  value);
+      break;
+    case STOP_SOLIB_GPU:
+      gdb_printf (file,
+		  _("Stopping for shared library events is \"%s\" "
+		    "(GPU code objects only).\n"),
+		  value);
+      break;
+    case STOP_SOLIB_ALL:
+      gdb_printf (file,
+		  _("Stopping for shared library events is \"%s\" "
+		    "(CPU and GPU).\n"),
+		  value);
+      break;
+    default:
+      gdb_printf (file,
+		  _("Stopping for shared library events is \"%s\".\n"),
+		  value);
+      break;
+    }
 }
 
 /* True after stop if current stack frame should be printed.  */
@@ -6462,7 +6506,8 @@ handle_inferior_event (struct execution_control_state *ecs)
 	       and place breakpoints in initializer routines for
 	       dynamically loaded objects (among other things).  */
 	    ecs->event_thread->set_stop_signal (GDB_SIGNAL_0);
-	    if (stop_on_solib_events)
+	    if (stop_on_solib_events == STOP_SOLIB_CPU
+		|| stop_on_solib_events == STOP_SOLIB_ALL)
 	      {
 		/* Make sure we print "Stopped due to solib-event" in
 		   normal_stop.  */
@@ -11032,13 +11077,18 @@ leave it stopped or free to run as needed."),
   /* Update cached state.  */
   signal_cache_update (-1);
 
-  add_setshow_zinteger_cmd ("stop-on-solib-events", class_support,
-			    &stop_on_solib_events, _("\
+  add_setshow_uinteger_cmd ("stop-on-solib-events", class_support,
+			    (unsigned int *) &stop_on_solib_events,
+			    stop_on_solib_events_literals, _("\
 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."),
+If a number, valid values are 0 through 3.  Alternatively, the following\n\
+keywords may be used:\n\
+  none - Do not stop on shared library events.\n\
+  cpu  - Stop on CPU shared library events only.\n\
+  gpu  - Stop on GPU code object events only.\n\
+  all  - Stop on both CPU and GPU library events.\n\
+The most common events of interest are loading/unloading of libraries."),
 			    set_stop_on_solib_events,
 			    show_stop_on_solib_events,
 			    &setlist, &showlist);

gdb/solib-svr4.c

@@ -2305,9 +2305,13 @@ svr4_update_solib_event_breakpoint (struct breakpoint *b)
 
       if (pa->action == DO_NOTHING)
 	{
-	  if (b->enable_state == bp_disabled && stop_on_solib_events)
+	  if (b->enable_state == bp_disabled
+	      && (stop_on_solib_events == STOP_SOLIB_CPU
+		  || stop_on_solib_events == STOP_SOLIB_ALL))
 	    enable_breakpoint (b);
-	  else if (b->enable_state == bp_enabled && !stop_on_solib_events)
+	  else if (b->enable_state == bp_enabled
+		   && stop_on_solib_events != STOP_SOLIB_CPU
+		   && stop_on_solib_events != STOP_SOLIB_ALL)
 	    disable_breakpoint (b);
 	}
 

gdb/solib.h

@@ -33,6 +33,13 @@ struct program_space;
 
 extern bool debug_solib;
 
+/* Values for 'set stop-on-solib-events' setting.  */
+
+#define STOP_SOLIB_NONE 0
+#define STOP_SOLIB_CPU  1
+#define STOP_SOLIB_GPU  2
+#define STOP_SOLIB_ALL  3
+
 /* Print an "solib" debug statement.  */
 
 #define solib_debug_printf(fmt, ...) \

gdb/testsuite/gdb.rocm/gpu-solib-event-kernel.cpp

@@ -0,0 +1,28 @@
+/* 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 "rocm-test-utils.h"
+#include <hip/hip_runtime.h>
+
+extern "C" __global__ void
+test_kernel ()
+{
+  NOP (1);
+}

gdb/testsuite/gdb.rocm/gpu-solib-event.cpp

@@ -0,0 +1,92 @@
+/* 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<char> module_buffer (module_size);
+
+  if (!mod.read (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));
+}
+
+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];
+
+  test_file_load (module_path);
+  test_memory_load (module_path);
+
+  return 0;
+}