diff --git a/gdb/amd-dbgapi-target.c b/gdb/amd-dbgapi-target.c index f3bab7e60cc..d2f9b531036 100644 --- a/gdb/amd-dbgapi-target.c +++ b/gdb/amd-dbgapi-target.c @@ -220,11 +220,13 @@ struct amd_dbgapi_inferior_info { explicit amd_dbgapi_inferior_info (inferior *inf, bool precise_memory_requested = false, - bool precise_alu_exceptions_requested = false) + bool precise_alu_exceptions_requested = false, + bool trap_on_entry_requested = false) : inf (inf) { precise_memory.requested = precise_memory_requested; precise_alu_exceptions.requested = precise_alu_exceptions_requested; + trap_on_entry.requested = trap_on_entry_requested; } /* Backlink to inferior. */ @@ -264,6 +266,9 @@ struct amd_dbgapi_inferior_info enablement. */ struct dbgapi_feature_request precise_alu_exceptions; + /* Track the status of trap-on-entry request and enablement. */ + struct dbgapi_feature_request trap_on_entry; + gdb::unordered_map breakpoint_map; @@ -2447,6 +2452,31 @@ set_process_alu_exceptions_precision (amd_dbgapi_inferior_info &info) get_status_string (status)); } +/* Set the process' wave launch trap-on-entry mode. */ + +static void +set_process_trap_on_entry (amd_dbgapi_inferior_info &info) +{ +#if HAVE_DECL_AMD_DBGAPI_PROCESS_ENABLE_TRAPS + auto value = (info.trap_on_entry.requested + ? AMD_DBGAPI_WAVE_ENABLE_TRAP_ON_ENTRY + : AMD_DBGAPI_WAVE_ENABLE_TRAP_NONE); + + amd_dbgapi_status_t status + = amd_dbgapi_process_enable_traps (info.process_id, value, AMD_DBGAPI_WAVE_ENABLE_TRAP_ON_ENTRY); +#else + amd_dbgapi_status_t status = AMD_DBGAPI_STATUS_ERROR_NOT_SUPPORTED; +#endif + + if (status == AMD_DBGAPI_STATUS_SUCCESS) + info.trap_on_entry.enabled = info.trap_on_entry.requested; + else if (status == AMD_DBGAPI_STATUS_ERROR_NOT_SUPPORTED) + warning (_("AMDGPU enabled traps could not be set.")); + else + error (_("amd_dbgapi_process_enable_traps failed (%s)"), + get_status_string (status)); +} + /* Handle extra initialisation after we have attached to a AMDGPU corefile. */ static void @@ -2570,6 +2600,7 @@ attach_amd_dbgapi (inferior *inf) set_process_memory_precision (info); set_process_alu_exceptions_precision (info); + set_process_trap_on_entry (info); /* If GDB is attaching to a process that has the runtime loaded, there will already be a "runtime loaded" event available. Consume it and push the @@ -2619,10 +2650,11 @@ detach_amd_dbgapi (inferior *inf) for (auto &&value : info.breakpoint_map) delete_breakpoint (value.second); - /* Reset the amd_dbgapi_inferior_info, except for precise_memory_mode and - precise_alu_exceptions. */ + /* Reset the amd_dbgapi_inferior_info, except for precise_memory_mode, + precise_alu_exceptions and trap_on_entry. */ info = amd_dbgapi_inferior_info (inf, info.precise_memory.requested, - info.precise_alu_exceptions.requested); + info.precise_alu_exceptions.requested, + info.trap_on_entry.requested); maybe_reset_amd_dbgapi (); } @@ -3058,6 +3090,7 @@ amd_dbgapi_target_inferior_cloned (inferior *original_inferior, new_info.precise_memory.requested = orig_info.precise_memory.requested; new_info.precise_alu_exceptions.requested = orig_info.precise_alu_exceptions.requested; + new_info.trap_on_entry.requested = orig_info.trap_on_entry.requested; } /* inferior_execd observer. */ @@ -3078,6 +3111,8 @@ amd_dbgapi_inferior_execd (inferior *exec_inf, inferior *follow_inf) get_amd_dbgapi_inferior_info (follow_inf).precise_alu_exceptions.requested = get_amd_dbgapi_inferior_info (exec_inf) .precise_alu_exceptions.requested; + get_amd_dbgapi_inferior_info (follow_inf).trap_on_entry.requested + = get_amd_dbgapi_inferior_info (exec_inf).trap_on_entry.requested; attach_amd_dbgapi (follow_inf); } @@ -3099,6 +3134,8 @@ amd_dbgapi_inferior_forked (inferior *parent_inf, inferior *child_inf, = parent_info.precise_memory.requested; child_info.precise_alu_exceptions.requested = parent_info.precise_alu_exceptions.requested; + child_info.trap_on_entry.requested + = parent_info.trap_on_entry.requested; if (fork_kind != TARGET_WAITKIND_VFORKED) { @@ -3663,6 +3700,55 @@ get_effective_precise_alu_exception_mode () return info.precise_alu_exceptions.enabled; } +/* Callback for "show amdgpu trap-on-entry". */ + +static void +show_trap_on_entry_mode (struct ui_file *file, int from_tty, + struct cmd_list_element *c, const char *value) +{ + const amd_dbgapi_inferior_info &info + = get_amd_dbgapi_inferior_info (current_inferior ()); + + gdb_printf (file, + _("AMDGPU wave launch trap-on-entry is %s (currently %s).\n"), + info.trap_on_entry.requested ? "on" : "off", + info.trap_on_entry.enabled ? "enabled" : "disabled"); +} + +/* Callback for "set amdgpu trap-on-entry". */ + +static void +set_trap_on_entry_mode (bool value) +{ + amd_dbgapi_inferior_info &info + = get_amd_dbgapi_inferior_info (current_inferior ()); + + info.trap_on_entry.requested = value; + + if (info.runtime_state == AMD_DBGAPI_RUNTIME_STATE_LOADED_SUCCESS) + set_process_trap_on_entry (info); +} + +/* Get the trap-on-entry requested mode. */ + +static bool +get_trap_on_entry_mode () +{ + const amd_dbgapi_inferior_info &info + = get_amd_dbgapi_inferior_info (current_inferior ()); + return info.trap_on_entry.requested; +} + +/* Get the trap-on-entry effective mode. */ + +static bool +get_effective_trap_on_entry_mode () +{ + amd_dbgapi_inferior_info &info + = get_amd_dbgapi_inferior_info (current_inferior ()); + return info.trap_on_entry.enabled; +} + static const char * get_dbgapi_library_file_path () { @@ -5014,6 +5100,19 @@ running. If off (default), precise ALU exceptions reporting is disabled."), cmds.show->var->set_effective_value_getter (get_effective_precise_alu_exception_mode); + cmds = add_setshow_boolean_cmd ("trap-on-entry", no_class, + _("Set trap-on-entry mode."), + _("Show trap-on-entry mode."), _("\ +If on, all newly created waves will trap immediately upon launch before\n\ +executing any shader instructions. If off (default), waves execute normally."), + set_trap_on_entry_mode, + get_trap_on_entry_mode, + show_trap_on_entry_mode, + &set_amdgpu_list, &show_amdgpu_list); + + cmds.show->var->set_effective_value_getter + (get_effective_trap_on_entry_mode); + add_cmd ("version", no_set_class, show_dbgapi_version, _("Show the ROCdbgapi library version and build information."), &show_amdgpu_list); diff --git a/gdb/config.in b/gdb/config.in index 73d0b1913ce..80205fbeac2 100644 --- a/gdb/config.in +++ b/gdb/config.in @@ -136,6 +136,11 @@ you don't. */ #undef HAVE_DECL_ADDR_NO_RANDOMIZE +/* Define to 1 if you have the declaration of + `amd_dbgapi_process_enable_traps', and to 0 if you don't. + */ +#undef HAVE_DECL_AMD_DBGAPI_PROCESS_ENABLE_TRAPS + /* Define to 1 if you have the declaration of `asprintf', and to 0 if you don't. */ #undef HAVE_DECL_ASPRINTF diff --git a/gdb/configure b/gdb/configure index 4674c85faac..8e8b3b52bd6 100755 --- a/gdb/configure +++ b/gdb/configure @@ -25437,6 +25437,22 @@ $as_echo "#define HAVE_AMD_DBGAPI 1" >>confdefs.h TARGET_OBS="$TARGET_OBS amd-dbgapi-target.o" + save_CPPFLAGS="$CPPFLAGS" + CPPFLAGS="$CPPFLAGS $AMD_DBGAPI_CFLAGS" + ac_fn_c_check_decl "$LINENO" "amd_dbgapi_process_enable_traps" "ac_cv_have_decl_amd_dbgapi_process_enable_traps" "#include +" +if test "x$ac_cv_have_decl_amd_dbgapi_process_enable_traps" = xyes; then : + ac_have_decl=1 +else + ac_have_decl=0 +fi + +cat >>confdefs.h <<_ACEOF +#define HAVE_DECL_AMD_DBGAPI_PROCESS_ENABLE_TRAPS $ac_have_decl +_ACEOF + + CPPFLAGS="$save_CPPFLAGS" + # If --enable-targets=all was provided, use the list of all files depending # on amd-dbgapi that is hardcoded in the Makefile. Else, the appropriate # architecture entry in configure.tgt will have added the files to diff --git a/gdb/configure.ac b/gdb/configure.ac index 28dd2ad14c7..b98da260cae 100644 --- a/gdb/configure.ac +++ b/gdb/configure.ac @@ -334,6 +334,12 @@ if test "$gdb_require_amd_dbgapi" = true \ AC_DEFINE(HAVE_AMD_DBGAPI, 1, [Define if amd-dbgapi is being linked in.]) TARGET_OBS="$TARGET_OBS amd-dbgapi-target.o" + save_CPPFLAGS="$CPPFLAGS" + CPPFLAGS="$CPPFLAGS $AMD_DBGAPI_CFLAGS" + AC_CHECK_DECLS([amd_dbgapi_process_enable_traps], [], [], + [[#include ]]) + CPPFLAGS="$save_CPPFLAGS" + # If --enable-targets=all was provided, use the list of all files depending # on amd-dbgapi that is hardcoded in the Makefile. Else, the appropriate # architecture entry in configure.tgt will have added the files to diff --git a/gdb/doc/gdb.texinfo b/gdb/doc/gdb.texinfo index b97e64aafef..0a08e28eb5e 100644 --- a/gdb/doc/gdb.texinfo +++ b/gdb/doc/gdb.texinfo @@ -3654,12 +3654,12 @@ Adds @var{n} inferiors ready to execute the same program as inferior @var{infno}; @var{n} defaults to 1, and @var{infno} defaults to the number of the current inferior. This command copies the values of the @var{args}, @w{@var{inferior-tty}}, @var{cwd}, @w{@code{amdgpu -precise-memory}} and @w{@code{amdgpu precise-alu-exceptions}} properties -from the current inferior to the new one. It also propagates changes -the user made to environment variables using the @w{@code{set -environment}} and @w{@code{unset environment}} commands. This is a -convenient command when you want to run another instance of the inferior -you are debugging. +precise-memory}}, @w{@code{amdgpu precise-alu-exceptions}} and +@w{@code{amdgpu trap-on-entry}} properties from the current inferior to +the new one. It also propagates changes the user made to environment +variables using the @w{@code{set environment}} and @w{@code{unset +environment}} commands. This is a convenient command when you want to +run another instance of the inferior you are debugging. @smallexample (@value{GDBP}) info inferiors @@ -30074,6 +30074,78 @@ Display the currently requested and effective @acronym{AMD} @end table +@subsubsection @acronym{AMD} @acronym{GPU} Forcing Waves to Trap on Entry + +By default, @acronym{AMD} @acronym{GPU} waves execute freely once they are +created, and are only stopped when they hit a breakpoint or report an +exception. It can be useful to instead have every wave stop as soon as it +enters a kernel, for example to inspect any activity on the device. + +The following commands can be used to control whether waves trap on entry: + +@table @code + +@kindex set amdgpu trap-on-entry +@cindex AMD GPU trap-on-entry +@item set amdgpu trap-on-entry @var{mode} +Controls whether @acronym{AMD} @acronym{GPU} waves trap when they enter a +kernel: + +@table @code + +@item off +The program executes normally. This is the default. + +@item on +Request that the program breaks at every kernel entry on the devices. + +The trap-on-entry mode can be changed at any time in the life of the +program. If the trap-on-entry mode is updated before the inferior +is started, the update request is recorded and applies as soon as +possible. For example: + +@smallexample +(@value{GDBP}) set amdgpu trap-on-entry on +(@value{GDBP}) show amdgpu trap-on-entry +AMDGPU wave launch trap-on-entry is on (currently disabled). +(@value{GDBP}) start +Temporary breakpoint 1 at 0x218cc9: file /tmp/trap-on-entry.cpp, line 46. +Starting program: /tmp/a.out +[Thread debugging using libthread_db enabled] +Using host libthread_db library "/lib/x86_64-linux-gnu/libthread_db.so.1". + +Temporary breakpoint 1, main (argc=1, argv=0x7fffffffe628) at /tmp/trap-on-entry.cpp:46 +46 if (argc != 2) +(@value{GDBP}) show amdgpu trap-on-entry +AMDGPU wave launch trap-on-entry is on (currently enabled). +@end smallexample + +Any activity on the GPU can thus be captured: + +@smallexample +(@value{GDBP}) continue +Continuing. + +Thread 15 "do_an_addition" received signal SIGTRAP, Trace/breakpoint trap. +[Switching to thread 15, lane 0 (AMDGPU Lane 1:2:2:9/0 (0,0,0)[0,0,0])] +do_an_addition () at /tmp/trap-on-entry.cpp:23 +23 @{ +@end smallexample + +This feature may expose some internal components from the runtime, which +may add clutter to the debugging session. It is advised to ensure these +are initialized prior to enabling it. + +@end table + +@kindex show amdgpu trap-on-entry +@cindex AMD GPU trap-on-entry +@item show amdgpu trap-on-entry +Display the currently requested and effective @acronym{AMD} +@acronym{GPU} trap-on-entry setting. + +@end table + @subsubsection @acronym{AMD} @acronym{GPU} Logging The @samp{set debug amd-dbgapi-lib log-level @var{level}} command can be used diff --git a/gdb/testsuite/gdb.rocm/trap-on-entry.cpp b/gdb/testsuite/gdb.rocm/trap-on-entry.cpp new file mode 100644 index 00000000000..16465f137db --- /dev/null +++ b/gdb/testsuite/gdb.rocm/trap-on-entry.cpp @@ -0,0 +1,72 @@ +/* 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 "hip/hip_runtime.h" +#include + +#define CHECK(cmd) \ + { \ + hipError_t error = cmd; \ + if (error != hipSuccess) \ + { \ + fprintf (stderr, "error: '%s'(%d) at %s:%d\n", \ + hipGetErrorString (error), error, __FILE__, __LINE__); \ + exit (EXIT_FAILURE); \ + } \ + } + +__global__ void +do_an_addition (int a, int b, int *out) +{ + *out = a + b; +} + +static void +run_addition (int *result_ptr) +{ + do_an_addition<<>> (1, 2, result_ptr); + + CHECK(hipDeviceSynchronize ()); +} + +static void +enable_trap_on_entry () +{ +} + +int +main () +{ + int *result_ptr; + + CHECK(hipMalloc (&result_ptr, sizeof (int))); + + /* Dry run: this first dispatch triggers internal kernels (heap + setup, etc.). */ + run_addition (result_ptr); + + /* The debugger enables trap-on-entry here. */ + enable_trap_on_entry (); + + /* This dispatch is expected to trap on entry. */ + run_addition (result_ptr); + + /* Re-run the kernel after disabling trap-on-entry. */ + run_addition (result_ptr); + + return 0; +} diff --git a/gdb/testsuite/gdb.rocm/trap-on-entry.exp b/gdb/testsuite/gdb.rocm/trap-on-entry.exp new file mode 100644 index 00000000000..9807b527a74 --- /dev/null +++ b/gdb/testsuite/gdb.rocm/trap-on-entry.exp @@ -0,0 +1,97 @@ +# 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 . + +# Exercise the `trap-on-entry` feature. Attempt to make the debugger break +# when a kernel runs, then verify that it disables cleanly. + +load_lib rocm.exp + +require allow_hip_tests + +standard_testfile .cpp + +if {[build_executable "failed to prepare" $testfile $srcfile {debug hip}]} { + return +} + +proc do_test {} { + clean_restart $::testfile + + with_rocm_gpu_lock { + if {![runto_main]} { + return + } + + # The feature must be disabled by default. + gdb_test "show amdgpu trap-on-entry" \ + "AMDGPU wave launch trap-on-entry is off \\(currently disabled\\)\\." \ + "trap-on-entry off by default" + + # Let the dry run complete so that all runtime-internal kernels (heap + # setup, etc.) have executed. Enabling trap-on-entry before that would + # trap on entry to those internal kernels instead of do_an_addition. + gdb_breakpoint "enable_trap_on_entry" + gdb_continue_to_breakpoint "after dry run" + + gdb_test_no_output "set amdgpu trap-on-entry on" "trap-on-entry enabled at runtime" + gdb_test "show amdgpu trap-on-entry" \ + "AMDGPU wave launch trap-on-entry is on \\(currently enabled\\)\\." \ + "trap-on-entry shown enabled" + + # With the feature enabled, the next dispatch traps on entry. + gdb_test "continue" \ + "Thread $::decimal \"do_an_addition\" received signal SIGTRAP, Trace/breakpoint trap.*" + + # Disable the feature while stopped in the trapped kernel. + gdb_test_no_output "set amdgpu trap-on-entry off" "trap-on-entry disabled at runtime" + gdb_test "show amdgpu trap-on-entry" \ + "AMDGPU wave launch trap-on-entry is off \\(currently disabled\\)\\." \ + "trap-on-entry shown disabled" + + # The remaining dispatch must not trap; the program runs to completion. + gdb_continue_to_end "" continue 1 + } +} + +# Not every architecture supports trap-on-entry. Request the feature before +# the process starts, then let the process attach apply it and report through +# "show" whether it actually became effective. +set trap_on_entry_supported false +with_test_prefix "check_trap_on_entry_support" { + clean_restart $::testfile + + with_rocm_gpu_lock { + gdb_test_no_output "set amdgpu trap-on-entry on" + if {![runto_main]} { + return + } + + gdb_test_multiple "show amdgpu trap-on-entry" "" { + -re -wrap "AMDGPU wave launch trap-on-entry is on \\(currently (enabled|disabled)\\).*" { + if {$expect_out(1,string) eq "enabled"} { + set trap_on_entry_supported true + } + pass $gdb_test_name + } + } + } +} + +if {!$trap_on_entry_supported} { + unsupported "trap-on-entry not supported" + return +} + +do_test