diff --git a/gdb/amd-dbgapi-target.c b/gdb/amd-dbgapi-target.c index f3bab7e60cc..c04990fd790 100644 --- a/gdb/amd-dbgapi-target.c +++ b/gdb/amd-dbgapi-target.c @@ -38,6 +38,7 @@ #include "gdbthread.h" #include "inf-loop.h" #include "inferior.h" +#include "infrun.h" #include "location.h" #include "objfiles.h" #include "observable.h" @@ -45,6 +46,7 @@ #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> 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_mapstop = 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; diff --git a/gdb/breakpoint.c b/gdb/breakpoint.c index f88680bf4f0..e42331b7f6b 100644 --- a/gdb/breakpoint.c +++ b/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; diff --git a/gdb/infrun.c b/gdb/infrun.c index 58c66d6191e..b6cd174cdc5 100644 --- a/gdb/infrun.c +++ b/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); diff --git a/gdb/solib-svr4.c b/gdb/solib-svr4.c index f2966a424a4..3e272369c48 100644 --- a/gdb/solib-svr4.c +++ b/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); } diff --git a/gdb/solib.h b/gdb/solib.h index 662d467b0cc..0ece66d0d6b 100644 --- a/gdb/solib.h +++ 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, ...) \ diff --git a/gdb/testsuite/gdb.rocm/gpu-solib-event-kernel.cpp b/gdb/testsuite/gdb.rocm/gpu-solib-event-kernel.cpp new file mode 100644 index 00000000000..2f9b4cf04b5 --- /dev/null +++ b/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 . */ + +/* Device kernel for solib-event test. + Compiled with --cuda-device-only to produce .co file. */ + +#include "rocm-test-utils.h" +#include + +extern "C" __global__ void +test_kernel () +{ + NOP (1); +} diff --git a/gdb/testsuite/gdb.rocm/gpu-solib-event.cpp b/gdb/testsuite/gdb.rocm/gpu-solib-event.cpp new file mode 100644 index 00000000000..8a0922d3c01 --- /dev/null +++ b/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 . */ + +#include "rocm-test-utils.h" +#include +#include +#include + +/* Test file:// events (hipModuleLoad). */ +static void +test_file_load (const char *module_path) +{ + hipModule_t module; + CHECK (hipModuleLoad (&module, module_path)); + + hipFunction_t function; + CHECK (hipModuleGetFunction (&function, module, "test_kernel")); + + CHECK (hipModuleLaunchKernel (function, 1, 1, 1, 1, 1, 1, + 0, nullptr, nullptr, nullptr)); + + CHECK (hipDeviceSynchronize ()); + CHECK (hipModuleUnload (module)); +} + +/* Test memory:// events (hipModuleLoadData). */ +static void +test_memory_load (const char *module_path) +{ + /* Read module file into memory buffer. */ + std::ifstream mod (module_path, std::ios::binary | std::ios::ate); + if (!mod.is_open ()) + { + fprintf (stderr, "Failed to open module file\n"); + exit (EXIT_FAILURE); + } + + size_t module_size = mod.tellg (); + mod.seekg (0, std::ios::beg); + std::vector module_buffer (module_size); + + if (!mod.read (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 \n", argv[0]); + return EXIT_FAILURE; + } + + const char *module_path = argv[1]; + + test_file_load (module_path); + test_memory_load (module_path); + + return 0; +} diff --git a/gdb/testsuite/gdb.rocm/gpu-solib-event.exp b/gdb/testsuite/gdb.rocm/gpu-solib-event.exp new file mode 100644 index 00000000000..18343b10912 --- /dev/null +++ b/gdb/testsuite/gdb.rocm/gpu-solib-event.exp @@ -0,0 +1,119 @@ +# 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 . + +# Test that "set stop-on-solib-events" works for GPU code objects. +# Tests both file:// events (hipModuleLoad) and memory:// events +# (hipModuleLoadData), as well as unload events (hipModuleUnload). +# +# Uses stop-on-solib-events value 2 (GPU events only) to avoid +# stopping on CPU shared library events. + +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 +} + +# Helper proc to continue until we see the expected GPU event. +# HIP runtime may load internal memory:// code objects, so we need to +# skip GPU events that don't match our specific criteria. +# Arguments: +# event_type: "load" or "unload" +# uri_pattern: regex pattern to match (e.g., "file://" or "memory://") +# test_name: descriptive name for the test +proc find_gpu_event {event_type uri_pattern test_name} { + set seen_gpu_event 0 + set max_continues 20 + + for {set i 0} {$i < $max_continues} {incr i} { + gdb_test_multiple "continue" "continue looking for $test_name" { + -re -wrap "Stopped due to GPU code object event.*Inferior (loaded|unloaded) ${uri_pattern}\[^\r\n\]+" { + if {[string equal $event_type "load"] && [string equal $expect_out(1,string) "loaded"]} { + set seen_gpu_event 1 + break + } elseif {[string equal $event_type "unload"] && [string equal $expect_out(1,string) "unloaded"]} { + set seen_gpu_event 1 + break + } + # GPU event doesn't match our criteria (wrong operation type), continue. + } + -re -wrap "Stopped due to GPU code object event.*" { + # GPU event with different URI type, continue. + } + -re -wrap "Inferior.*exited" { + break + } + -re -wrap ".*" { + fail "unexpected stop before $test_name" + break + } + } + + if {$seen_gpu_event} { + break + } + } + + gdb_assert {$seen_gpu_event} $test_name +} + +proc do_test {} { + with_rocm_gpu_lock { + clean_restart + gdb_load $::binfile + + if {![runto_main -inferior-args $::hipmodule_path]} { + return + } + + # Use "gpu" to stop on GPU code object events only. + # This avoids CPU shared library events during test execution. + gdb_test_no_output "set stop-on-solib-events gpu" + + # Test 1: file:// load event (hipModuleLoad). + gdb_breakpoint "test_file_load" -temporary + gdb_continue_to_breakpoint "at test_file_load" + find_gpu_event "load" "file://" "saw GPU code object load event with file:// prefix" + + # Test 2: file:// unload event (hipModuleUnload). + find_gpu_event "unload" "file://" "saw GPU code object unload event with file:// prefix" + + # Test 3: memory:// load event (hipModuleLoadData). + gdb_breakpoint "test_memory_load" -temporary + gdb_continue_to_breakpoint "at test_memory_load" + find_gpu_event "load" "memory://" "saw GPU code object load event with memory:// prefix" + + # Test 4: memory:// unload event (hipModuleUnload). + find_gpu_event "unload" "memory://" "saw GPU code object unload event with memory:// prefix" + } +} + +do_test