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

[Lancelot SIX <lancelot.six@amd.com>]

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