* [PATCH 1/2] gdb: add inferior_cloned observable @ 2023-09-13 15:28 Simon Marchi 2023-09-13 15:28 ` [PATCH 2/2] gdb/amdgpu: add precise-memory support Simon Marchi 2023-09-14 14:08 ` [PATCH 1/2] gdb: add inferior_cloned observable Tom Tromey 0 siblings, 2 replies; 10+ messages in thread From: Simon Marchi @ 2023-09-13 15:28 UTC (permalink / raw) To: gdb-patches; +Cc: Simon Marchi The following patch makes the amdgpu port transfer a property from the original inferior to the new inferior when using the clone-inferior command. Add the inferior_cloned observable to help with this. Change-Id: Id845a799813ec49b1b7b2fcb97b07d0a1e5e2631 --- gdb/inferior.c | 2 ++ gdb/observable.c | 1 + gdb/observable.h | 4 ++++ 3 files changed, 7 insertions(+) diff --git a/gdb/inferior.c b/gdb/inferior.c index ce4960a508aa..550bbd2827c0 100644 --- a/gdb/inferior.c +++ b/gdb/inferior.c @@ -1041,6 +1041,8 @@ clone_inferior_command (const char *args, int from_tty) for (const std::string &unset_var : orginf->environment.user_unset_env ()) inf->environment.unset (unset_var.c_str ()); + + gdb::observers::inferior_cloned.notify (orginf, inf); } } diff --git a/gdb/observable.c b/gdb/observable.c index a18b863fb471..33e51223cf41 100644 --- a/gdb/observable.c +++ b/gdb/observable.c @@ -58,6 +58,7 @@ DEFINE_OBSERVABLE (inferior_appeared); DEFINE_OBSERVABLE (inferior_pre_detach); DEFINE_OBSERVABLE (inferior_exit); DEFINE_OBSERVABLE (inferior_removed); +DEFINE_OBSERVABLE (inferior_cloned); DEFINE_OBSERVABLE (memory_changed); DEFINE_OBSERVABLE (before_prompt); DEFINE_OBSERVABLE (gdb_datadir_changed); diff --git a/gdb/observable.h b/gdb/observable.h index c0bafc51f143..4ea203c6fc38 100644 --- a/gdb/observable.h +++ b/gdb/observable.h @@ -172,6 +172,10 @@ extern observable<struct inferior */* inf */> inferior_exit; This method is called immediately before freeing INF. */ extern observable<struct inferior */* inf */> inferior_removed; +/* The inferior CLONE has been created by cloning INF. */ +extern observable<struct inferior */* inf */, struct inferior */* clone */> + inferior_cloned; + /* Bytes from DATA to DATA + LEN have been written to the inferior at ADDR. */ extern observable<struct inferior */* inferior */, CORE_ADDR /* addr */, base-commit: 38cc67cc00dc7385991e2658a1b4213d5ba3c4f0 -- 2.42.0 ^ permalink raw reply [flat|nested] 10+ messages in thread
* [PATCH 2/2] gdb/amdgpu: add precise-memory support 2023-09-13 15:28 [PATCH 1/2] gdb: add inferior_cloned observable Simon Marchi @ 2023-09-13 15:28 ` Simon Marchi 2023-09-13 21:32 ` Lancelot SIX 2023-09-14 14:10 ` Tom Tromey 2023-09-14 14:08 ` [PATCH 1/2] gdb: add inferior_cloned observable Tom Tromey 1 sibling, 2 replies; 10+ messages in thread From: Simon Marchi @ 2023-09-13 15:28 UTC (permalink / raw) To: gdb-patches; +Cc: Simon Marchi The amd-dbgapi library exposes a setting called "memory precision" for AMD GPUs [1]. Here's a copy of the description of the setting: The AMD GPU can overlap the execution of memory instructions with other instructions. This can result in a wave stopping due to a memory violation or hardware data watchpoint hit with a program counter beyond the instruction that caused the wave to stop. Some architectures allow the hardware to be configured to always wait for memory operations to complete before continuing. This will result in the wave stopping at the instruction immediately after the one that caused the stop event. Enabling this mode can make execution of waves significantly slower. Expose this option through a new "amdgpu precise-memory" setting. The precise memory setting is per inferior. The setting is transferred from one inferior to another when using the clone-inferior command, or when a new inferior is created following an exec or a fork. It can be set before starting the inferior, in which case GDB will attempt to apply what the user wants when attaching amd-dbgapi. If the user has requested to enable precise memory, but it can't be enabled (not all hardware supports it), GDB prints a warning. If precise memory is disabled, GDB prints a warning when hitting a memory exception (translated into GDB_SIGNAL_SEGV or GDB_SIGNAL_BUS), saying that the stop location may not be precise. Note that the precise memory setting also affects memory watchpoint reporting, but the watchpoint support for AMD GPUs hasn't been upstreamed to GDB yet. When we do upstream watchpoint support, GDB will produce a similar warning message when stopping due to a watchpoint if precise memory is disabled. Add a handful of tests. Add a util proc "hip_device_supports_precise_memory", which indicates if the device used for testing supports that feature. To implement it, also add a new "hcc_amdgpu_target" proc, to return the architecture of the device used for testing. [1] https://github.com/ROCm-Developer-Tools/ROCdbgapi/blob/687374258a27b5aab1309a7e8ded719e2f1ed3b1/include/amd-dbgapi.h.in#L6300-L6317 Change-Id: Ife1a99c0e960513da375ced8f8afaf8e47a61b3f --- gdb/amd-dbgapi-target.c | 199 +++++++++++++++++- gdb/doc/gdb.texinfo | 43 ++++ gdb/testsuite/gdb.rocm/precise-memory-exec.c | 44 ++++ .../gdb.rocm/precise-memory-exec.exp | 62 ++++++ gdb/testsuite/gdb.rocm/precise-memory-fork.c | 41 ++++ .../gdb.rocm/precise-memory-fork.exp | 54 +++++ .../precise-memory-multi-inferiors.exp | 87 ++++++++ .../precise-memory-warning-sigsegv.cpp | 33 +++ .../precise-memory-warning-sigsegv.exp | 49 +++++ gdb/testsuite/gdb.rocm/precise-memory.cpp | 32 +++ gdb/testsuite/gdb.rocm/precise-memory.exp | 57 +++++ gdb/testsuite/lib/rocm.exp | 59 ++++++ 12 files changed, 755 insertions(+), 5 deletions(-) create mode 100644 gdb/testsuite/gdb.rocm/precise-memory-exec.c create mode 100644 gdb/testsuite/gdb.rocm/precise-memory-exec.exp create mode 100644 gdb/testsuite/gdb.rocm/precise-memory-fork.c create mode 100644 gdb/testsuite/gdb.rocm/precise-memory-fork.exp create mode 100644 gdb/testsuite/gdb.rocm/precise-memory-multi-inferiors.exp create mode 100644 gdb/testsuite/gdb.rocm/precise-memory-warning-sigsegv.cpp create mode 100644 gdb/testsuite/gdb.rocm/precise-memory-warning-sigsegv.exp create mode 100644 gdb/testsuite/gdb.rocm/precise-memory.cpp create mode 100644 gdb/testsuite/gdb.rocm/precise-memory.exp diff --git a/gdb/amd-dbgapi-target.c b/gdb/amd-dbgapi-target.c index 22c269b7992c..cfb935df1163 100644 --- a/gdb/amd-dbgapi-target.c +++ b/gdb/amd-dbgapi-target.c @@ -23,6 +23,7 @@ #include "amdgpu-tdep.h" #include "async-event.h" #include "cli/cli-cmds.h" +#include "cli/cli-decode.h" #include "cli/cli-style.h" #include "inf-loop.h" #include "inferior.h" @@ -139,6 +140,17 @@ struct amd_dbgapi_inferior_info Initialized to true, since that's the default in amd-dbgapi too. */ bool forward_progress_required = true; + struct + { + /* Whether precise memory reporting is requested. */ + bool requested = false; + + /* Whether precise memory was requested and successfully enabled by + dbgapi (it may not be available for the current hardware, for + instance). */ + bool enabled = false; + } precise_memory; + std::unordered_map<decltype (amd_dbgapi_breakpoint_id_t::handle), struct breakpoint *> breakpoint_map; @@ -1326,6 +1338,36 @@ amd_dbgapi_target::stopped_by_hw_breakpoint () return false; } +/* Set the process's memory access reporting precision. + + The precision can be ::AMD_DBGAPI_MEMORY_PRECISION_PRECISE (waves wait for + memory instructions to complete before executing further instructions), or + ::AMD_DBGAPI_MEMORY_PRECISION_NONE (memory instructions execute normally). + + Returns true if the precision is supported by the architecture of all agents + in the process, or false if at least one agent does not support the + requested precision. + + An error is thrown if setting the precision results in a status other than + ::AMD_DBGAPI_STATUS_SUCCESS or ::AMD_DBGAPI_STATUS_ERROR_NOT_SUPPORTED. */ + +static bool +set_process_memory_precision (amd_dbgapi_process_id_t process_id, + amd_dbgapi_memory_precision_t precision) +{ + amd_dbgapi_status_t status + = amd_dbgapi_set_memory_precision (process_id, precision); + + if (status == AMD_DBGAPI_STATUS_ERROR_NOT_SUPPORTED) + return false; + + if (status != AMD_DBGAPI_STATUS_SUCCESS) + error (_("amd_dbgapi_set_memory_precision failed (%s)"), + get_status_string (status)); + + return true; +} + /* Make the amd-dbgapi library attach to the process behind INF. Note that this is unrelated to the "attach" GDB concept / command. @@ -1399,6 +1441,16 @@ attach_amd_dbgapi (inferior *inf) amd_dbgapi_debug_printf ("process_id = %" PRIu64 ", notifier fd = %d", info->process_id.handle, info->notifier); + amd_dbgapi_memory_precision_t memory_precision + = (info->precise_memory.requested + ? AMD_DBGAPI_MEMORY_PRECISION_PRECISE + : AMD_DBGAPI_MEMORY_PRECISION_NONE); + if (set_process_memory_precision (info->process_id, memory_precision)) + info->precise_memory.enabled = info->precise_memory.requested; + else + warning + (_("AMDGPU precise memory access reporting could not be enabled.")); + /* 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 target. */ @@ -1443,8 +1495,10 @@ detach_amd_dbgapi (inferior *inf) for (auto &&value : info->breakpoint_map) delete_breakpoint (value.second); - /* Reset the amd_dbgapi_inferior_info. */ + /* Reset the amd_dbgapi_inferior_info, except for precise_memory_mode. */ + bool precise_memory_requested = info->precise_memory.requested; *info = amd_dbgapi_inferior_info (inf); + info->precise_memory.requested = precise_memory_requested; maybe_reset_amd_dbgapi (); } @@ -1668,6 +1722,22 @@ amd_dbgapi_target_inferior_created (inferior *inf) attach_amd_dbgapi (inf); } +/* Callback called when an inferior is cloned. */ + +static void +amd_dbgapi_target_inferior_cloned (inferior *original_inferior, + inferior *new_inferior) +{ + auto *orig_info = get_amd_dbgapi_inferior_info (original_inferior); + auto *new_info = get_amd_dbgapi_inferior_info (new_inferior); + + /* At this point, the process is not started. Therefore it is sufficient to + copy the precise memory request, it will be applied when the process + starts. */ + gdb_assert (new_info->process_id == AMD_DBGAPI_PROCESS_NONE); + new_info->precise_memory.requested = orig_info->precise_memory.requested; +} + /* inferior_execd observer. */ static void @@ -1677,6 +1747,13 @@ amd_dbgapi_inferior_execd (inferior *exec_inf, inferior *follow_inf) attached to the old process image, so we need to detach and re-attach to the new process image. */ detach_amd_dbgapi (exec_inf); + + /* If using "follow-exec-mode new", carry over the precise-memory setting + to the new inferior (otherwise, FOLLOW_INF and ORIG_INF point to the same + inferior, so this is a no-op). */ + get_amd_dbgapi_inferior_info (follow_inf)->precise_memory.requested + = get_amd_dbgapi_inferior_info (exec_inf)->precise_memory.requested; + attach_amd_dbgapi (follow_inf); } @@ -1686,11 +1763,22 @@ static void amd_dbgapi_inferior_forked (inferior *parent_inf, inferior *child_inf, target_waitkind fork_kind) { - if (child_inf != nullptr && fork_kind != TARGET_WAITKIND_VFORKED) + if (child_inf != nullptr) { - scoped_restore_current_thread restore_thread; - switch_to_thread (*child_inf->threads ().begin ()); - attach_amd_dbgapi (child_inf); + /* Copy precise-memory requested value from parent to child. */ + amd_dbgapi_inferior_info *parent_info + = get_amd_dbgapi_inferior_info (parent_inf); + amd_dbgapi_inferior_info *child_info + = get_amd_dbgapi_inferior_info (child_inf); + child_info->precise_memory.requested + = parent_info->precise_memory.requested; + + if (fork_kind != TARGET_WAITKIND_VFORKED) + { + scoped_restore_current_thread restore_thread; + switch_to_thread (*child_inf->threads ().begin ()); + attach_amd_dbgapi (child_inf); + } } } @@ -1785,6 +1873,29 @@ amd_dbgapi_remove_breakpoint_callback return AMD_DBGAPI_STATUS_SUCCESS; } +/* signal_received observer. */ + +static void +amd_dbgapi_target_signal_received (gdb_signal sig) +{ + amd_dbgapi_inferior_info *info + = get_amd_dbgapi_inferior_info (current_inferior ()); + + if (info->process_id == AMD_DBGAPI_PROCESS_NONE) + return; + + if (!ptid_is_gpu (inferior_thread ()->ptid)) + return; + + if (sig != GDB_SIGNAL_SEGV && sig != GDB_SIGNAL_BUS) + return; + + if (!info->precise_memory.enabled) + gdb_printf ("\ +Warning: precise memory violation signal reporting is not enabled, reported\n\ +location may not be accurate. See \"show amdgpu precise-memory\".\n"); +} + /* Style for some kinds of messages. */ static cli_style_option fatal_error_style @@ -1853,6 +1964,62 @@ amd_dbgapi_target::close () delete_async_event_handler (&amd_dbgapi_async_event_handler); } +/* Callback for "show amdgpu precise-memory". */ + +static void +show_precise_memory_mode (struct ui_file *file, int from_tty, + struct cmd_list_element *c, const char *value) +{ + amd_dbgapi_inferior_info *info + = get_amd_dbgapi_inferior_info (current_inferior ()); + + gdb_printf (file, + _("AMDGPU precise memory access reporting is %s " + "(currently %s).\n"), + info->precise_memory.requested ? "on" : "off", + info->precise_memory.enabled ? "enabled" : "disabled"); +} + +/* Callback for "set amdgpu precise-memory". */ + +static void +set_precise_memory_mode (bool value) +{ + amd_dbgapi_inferior_info *info + = get_amd_dbgapi_inferior_info (current_inferior ()); + + info->precise_memory.requested = value; + + if (info->process_id != AMD_DBGAPI_PROCESS_NONE) + { + amd_dbgapi_memory_precision_t memory_precision + = (info->precise_memory.requested + ? AMD_DBGAPI_MEMORY_PRECISION_PRECISE + : AMD_DBGAPI_MEMORY_PRECISION_NONE); + + if (set_process_memory_precision (info->process_id, memory_precision)) + info->precise_memory.enabled = info->precise_memory.requested; + else + warning + (_("AMDGPU precise memory access reporting could not be enabled.")); + } +} + +/* Return whether precise-memory is requested for the current inferior. */ + +static bool +get_precise_memory_mode () +{ + amd_dbgapi_inferior_info *info + = get_amd_dbgapi_inferior_info (current_inferior ()); + + return info->precise_memory.requested; +} + +/* List of set/show amdgpu commands. */ +struct cmd_list_element *set_amdgpu_list; +struct cmd_list_element *show_amdgpu_list; + /* List of set/show debug amd-dbgapi-lib commands. */ struct cmd_list_element *set_debug_amd_dbgapi_lib_list; struct cmd_list_element *show_debug_amd_dbgapi_lib_list; @@ -1960,6 +2127,10 @@ _initialize_amd_dbgapi_target () amd_dbgapi_set_log_level (get_debug_amd_dbgapi_lib_log_level ()); /* Install observers. */ + gdb::observers::inferior_cloned.attach (amd_dbgapi_target_inferior_cloned, + "amd-dbgapi"); + gdb::observers::signal_received.attach (amd_dbgapi_target_signal_received, + "amd-dbgapi"); gdb::observers::inferior_created.attach (amd_dbgapi_target_inferior_created, amd_dbgapi_target_inferior_created_observer_token, "amd-dbgapi"); @@ -1968,6 +2139,24 @@ _initialize_amd_dbgapi_target () gdb::observers::inferior_exit.attach (amd_dbgapi_inferior_exited, "amd-dbgapi"); gdb::observers::inferior_pre_detach.attach (amd_dbgapi_inferior_pre_detach, "amd-dbgapi"); + add_basic_prefix_cmd ("amdgpu", no_class, + _("Generic command for setting amdgpu flags."), + &set_amdgpu_list, 0, &setlist); + + add_show_prefix_cmd ("amdgpu", no_class, + _("Generic command for showing amdgpu flags."), + &show_amdgpu_list, 0, &showlist); + + add_setshow_boolean_cmd ("precise-memory", no_class, + _("Set precise-memory mode."), + _("Show precise-memory mode."), _("\ +If on, precise memory reporting is enabled if/when the inferior is running.\n\ +If off (default), precise memory reporting is disabled."), + set_precise_memory_mode, + get_precise_memory_mode, + show_precise_memory_mode, + &set_amdgpu_list, &show_amdgpu_list); + add_basic_prefix_cmd ("amd-dbgapi-lib", no_class, _("Generic command for setting amd-dbgapi library " "debugging flags."), diff --git a/gdb/doc/gdb.texinfo b/gdb/doc/gdb.texinfo index 9b7e06f31566..fa91d72695e7 100644 --- a/gdb/doc/gdb.texinfo +++ b/gdb/doc/gdb.texinfo @@ -26794,6 +26794,49 @@ either not mapped or accessed with incompatible permissions. If a single instruction raises more than one signal, they will be reported one at a time each time the wavefront is continued. +@subsubsection @acronym{AMD GPU} Memory Violation Reporting + +A wavefront can report memory violation events. However, the program +location at which they are reported may be after the machine instruction +that caused them. This can result in the reported source statement +being incorrect. The following commands can be used to control this +behavior: + +@table @code + +@kindex set amdgpu precise-memory +@cindex AMD GPU precise memory event reporting +@item set amdgpu precise-memory @var{mode} +Controls how @acronym{AMD GPU} devices detect memory violations, where +@var{mode} can be: + +@table @code + +@item off +The program location may not be immediately after the instruction that +caused the memory violation. This is the default. + +@item on +Requests that the program location will be immediately after the +instruction that caused a memory violation. Enabling this mode may make +the @acronym{AMD GPU} device execution significantly slower as it has to +wait for each memory operation to complete before executing the next +instruction. + +@end table + +The @code{set amdgpu precise-memory} parameter is per-inferior. When an +inferior forks or execs, or the user uses the @code{clone-inferior} command, +and an inferior is created as a result, the newly created inferior inherits +the parameter value of the original inferior. + +@kindex show amdgpu precise-memory +@cindex AMD GPU precise memory event reporting +@item show amdgpu precise-memory +Displays the currently requested AMD GPU precise memory setting. + +@end table + @subsubsection @acronym{AMD GPU} Logging The @samp{set debug amd-dbgapi} command can be used diff --git a/gdb/testsuite/gdb.rocm/precise-memory-exec.c b/gdb/testsuite/gdb.rocm/precise-memory-exec.c new file mode 100644 index 000000000000..f0659a63fc5a --- /dev/null +++ b/gdb/testsuite/gdb.rocm/precise-memory-exec.c @@ -0,0 +1,44 @@ +/* Copyright 2021-2023 Free Software Foundation, Inc. + + This file is part of GDB. + + 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 <unistd.h> +#include <stdlib.h> +#include <stdio.h> + +static void +second (void) +{ +} + +int +main (int argc, char **argv) +{ + if (argc == 1) + { + /* First invocation */ + int ret = execl (argv[0], argv[0], "Hello", NULL); + perror ("exec"); + abort (); + } + else + { + /* Second invocation */ + second (); + } + + return 0; +} diff --git a/gdb/testsuite/gdb.rocm/precise-memory-exec.exp b/gdb/testsuite/gdb.rocm/precise-memory-exec.exp new file mode 100644 index 000000000000..26be6cf72146 --- /dev/null +++ b/gdb/testsuite/gdb.rocm/precise-memory-exec.exp @@ -0,0 +1,62 @@ +# Copyright 2021-2023 Free Software Foundation, Inc. + +# This file is part of GDB. + +# 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 the "set amdgpu precise-memory" setting is inherited by an inferior +# created following an exec. + +load_lib rocm.exp + +require allow_hipcc_tests + +if { ![istarget "*-linux*"] } then { + continue +} + +standard_testfile .c + +if {[build_executable "failed to prepare $testfile" $testfile $srcfile {debug}]} { + return +} + +proc do_test { follow-exec-mode } { + clean_restart $::binfile + + with_rocm_gpu_lock { + if ![runto_main] { + return + } + + # Set precise-memory on the inferior before exec. + gdb_test "show amdgpu precise-memory" " is off.*" \ + "show amdgpu precise-memory before set" + gdb_test "set amdgpu precise-memory on" + gdb_test "show amdgpu precise-memory" " is on.*" \ + "show amdgpu precise-memory after set" + + # Continue past exec. The precise-memory setting should + # be on. + gdb_test_no_output "set follow-exec-mode ${follow-exec-mode}" + gdb_test "break second" + gdb_test "continue" "Breakpoint 1(\.$::decimal)?, main .*" + gdb_test "show amdgpu precise-memory" " is on.*" \ + "show amdgpu precise-memory after exec" + } +} + +foreach_with_prefix follow-exec-mode {same new} { + do_test ${follow-exec-mode} +} diff --git a/gdb/testsuite/gdb.rocm/precise-memory-fork.c b/gdb/testsuite/gdb.rocm/precise-memory-fork.c new file mode 100644 index 000000000000..67ce09f2c3dc --- /dev/null +++ b/gdb/testsuite/gdb.rocm/precise-memory-fork.c @@ -0,0 +1,41 @@ +/* Copyright 2021-2023 Free Software Foundation, Inc. + + This file is part of GDB. + + 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 <unistd.h> + +static void +parent (void) +{ +} + +static void +child (void) +{ +} + +int +main (void) +{ + int pid = fork (); + + if (pid != 0) + parent (); + else + child (); + + return 0; +} diff --git a/gdb/testsuite/gdb.rocm/precise-memory-fork.exp b/gdb/testsuite/gdb.rocm/precise-memory-fork.exp new file mode 100644 index 000000000000..0dc88b89f8a7 --- /dev/null +++ b/gdb/testsuite/gdb.rocm/precise-memory-fork.exp @@ -0,0 +1,54 @@ +# Copyright 2021-2023 Free Software Foundation, Inc. + +# This file is part of GDB. + +# 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 the "set amdgpu precise-memory" setting is inherited by a fork +# child. + +load_lib rocm.exp + +require allow_hipcc_tests + +if { ![istarget "*-linux*"] } then { + continue +} + +standard_testfile .c + +if {[prepare_for_testing "failed to prepare $testfile" $testfile $srcfile {debug}]} { + return +} + +with_rocm_gpu_lock { + if ![runto_main] { + return + } + + # Set precise-memory on in the parent, before fork. + gdb_test "show amdgpu precise-memory" " is off.*" \ + "show amdgpu precise-memory before set" + gdb_test "set amdgpu precise-memory on" + gdb_test "show amdgpu precise-memory" " is on.*" \ + "show amdgpu precise-memory after set" + + # Continue past fork, following the child. The precise-memory setting should + # be on. + gdb_test "set follow-fork-mode child" + gdb_test "break child" + gdb_test "continue" "Thread 2.1 .* hit Breakpoint .*" + gdb_test "show amdgpu precise-memory" " is on.*" \ + "show amdgpu precise-memory after fork" +} diff --git a/gdb/testsuite/gdb.rocm/precise-memory-multi-inferiors.exp b/gdb/testsuite/gdb.rocm/precise-memory-multi-inferiors.exp new file mode 100644 index 000000000000..9968b422b0ee --- /dev/null +++ b/gdb/testsuite/gdb.rocm/precise-memory-multi-inferiors.exp @@ -0,0 +1,87 @@ +# Copyright 2021-2023 Free Software Foundation, Inc. + +# This file is part of GDB. + +# 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 the "set amdgpu precise-memory" setting is per-inferior, and +# inherited by an inferior created using the clone-inferior command. + +load_lib rocm.exp + +require allow_hipcc_tests + +clean_restart + +set test_python [allow_python_tests] + +proc test_per_inferior { } { + gdb_test "show amdgpu precise-memory" \ + "AMDGPU precise memory access reporting is off \\(currently disabled\\)." \ + "show initial value, inferior 1" + if $::test_python { + gdb_test "python print(gdb.parameter(\"amdgpu precise-memory\"))" \ + "False" \ + "show initial value using Python, inferior 1" + } + gdb_test_no_output "set amdgpu precise-memory" \ + "set on inferior 1" + gdb_test "show amdgpu precise-memory" \ + "AMDGPU precise memory access reporting is on \\(currently disabled\\)." \ + "show new value, inferior 1" + if $::test_python { + gdb_test "python print(gdb.parameter(\"amdgpu precise-memory\"))" \ + "True" \ + "show new value using Python, inferior 1" + } + + gdb_test "add-inferior" "Added inferior 2" + gdb_test "inferior 2" "Switching to inferior 2 .*" + + gdb_test "show amdgpu precise-memory" \ + "AMDGPU precise memory access reporting is off \\(currently disabled\\)." \ + "show initial value, inferior 2" + if $::test_python { + gdb_test "python print(gdb.parameter(\"amdgpu precise-memory\"))" \ + "False" \ + "show initial value using Python, inferior 2" + } +} + +proc test_copy_precise_memory_on_clone {precise_memory} { + set value $precise_memory + if {$precise_memory == "unspecified"} { + set value off + } + + clean_restart + gdb_test "show amdgpu precise-memory" "is off.*" \ + "show default amdgpu precise-memory" + if {$precise_memory != "unspecified"} { + gdb_test_no_output "set amdgpu precise-memory $value" + gdb_test "show amdgpu precise-memory" "is $value.*" \ + "show amdgpu precise-memory on original inferior" + } + + gdb_test "clone-inferior" "Added inferior 2.*" + gdb_test "inferior 2" + gdb_test "show amdgpu precise-memory" "is $value.*" \ + "show amdgpu precise-memory on cloned inferior" +} + +test_per_inferior + +foreach_with_prefix precise_memory { unspecified on off } { + test_copy_precise_memory_on_clone $precise_memory +} diff --git a/gdb/testsuite/gdb.rocm/precise-memory-warning-sigsegv.cpp b/gdb/testsuite/gdb.rocm/precise-memory-warning-sigsegv.cpp new file mode 100644 index 000000000000..58339e5391a6 --- /dev/null +++ b/gdb/testsuite/gdb.rocm/precise-memory-warning-sigsegv.cpp @@ -0,0 +1,33 @@ +/* Copyright 2021-2023 Free Software Foundation, Inc. + + This file is part of GDB. + + 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 <hip/hip_runtime.h> + +__global__ void +kernel () +{ + int *p = nullptr; + *p = 1; +} + +int +main (int argc, char* argv[]) +{ + hipLaunchKernelGGL (kernel, dim3 (1), dim3 (1), 0, 0); + hipDeviceSynchronize (); + return 0; +} diff --git a/gdb/testsuite/gdb.rocm/precise-memory-warning-sigsegv.exp b/gdb/testsuite/gdb.rocm/precise-memory-warning-sigsegv.exp new file mode 100644 index 000000000000..22e1f6eda254 --- /dev/null +++ b/gdb/testsuite/gdb.rocm/precise-memory-warning-sigsegv.exp @@ -0,0 +1,49 @@ +# Copyright 2021-2023 Free Software Foundation, Inc. + +# This file is part of GDB. + +# 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 when "amdgpu precise-memory" is off, hitting a SIGSEGV shows a +# warning about the stop location maybe being inaccurate. + +load_lib rocm.exp + +require allow_hipcc_tests + +if { ![istarget "*-linux*"] } then { + continue +} + +standard_testfile .cpp + +if {[build_executable "failed to prepare" $testfile $srcfile {debug hip}]} { + return +} + +proc do_test { } { + clean_restart $::binfile + + with_rocm_gpu_lock { + if ![runto_main] { + return + } + + gdb_test_no_output "set amdgpu precise-memory off" + gdb_test "continue" \ + "SIGSEGV, Segmentation fault.*Warning: precise memory violation signal reporting is not enabled.*" + } +} + +do_test diff --git a/gdb/testsuite/gdb.rocm/precise-memory.cpp b/gdb/testsuite/gdb.rocm/precise-memory.cpp new file mode 100644 index 000000000000..6e0a4d9bc63b --- /dev/null +++ b/gdb/testsuite/gdb.rocm/precise-memory.cpp @@ -0,0 +1,32 @@ +/* Copyright 2021-2023 Free Software Foundation, Inc. + + This file is part of GDB. + + 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 <hip/hip_runtime.h> + +__global__ void +kernel () +{ + __builtin_amdgcn_s_sleep (1); +} + +int +main (int argc, char* argv[]) +{ + hipLaunchKernelGGL (kernel, dim3 (1), dim3 (1), 0, 0); + hipDeviceSynchronize (); + return 0; +} diff --git a/gdb/testsuite/gdb.rocm/precise-memory.exp b/gdb/testsuite/gdb.rocm/precise-memory.exp new file mode 100644 index 000000000000..bd2b12747c6f --- /dev/null +++ b/gdb/testsuite/gdb.rocm/precise-memory.exp @@ -0,0 +1,57 @@ +# Copyright 2022-2023 Free Software Foundation, Inc. + +# This file is part of GDB. + +# 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 showing the "amdgpu precise-memory" setting. + +load_lib rocm.exp + +require allow_hipcc_tests + +standard_testfile .cpp + +if {[build_executable "failed to prepare" $testfile $srcfile {debug hip}]} { + return +} + +proc do_test { } { + clean_restart $::binfile + + with_rocm_gpu_lock { + if ![runto_main] { + return + } + + gdb_test "show amdgpu precise-memory" \ + "AMDGPU precise memory access reporting is off \\(currently disabled\\)." \ + "show precise-memory setting in CLI before" + + if {[hip_device_supports_precise_memory]} { + gdb_test_no_output "set amdgpu precise-memory on" + set cli_effective_value "enabled" + } else { + gdb_test "set amdgpu precise-memory on" \ + "warning: AMDGPU precise memory access reporting could not be enabled." + set cli_effective_value "disabled" + } + + gdb_test "show amdgpu precise-memory" \ + "AMDGPU precise memory access reporting is on \\(currently ${cli_effective_value}\\)." \ + "show precise-memory setting in CLI after" + } +} + +do_test diff --git a/gdb/testsuite/lib/rocm.exp b/gdb/testsuite/lib/rocm.exp index 98a3b308228d..22b294a5efae 100644 --- a/gdb/testsuite/lib/rocm.exp +++ b/gdb/testsuite/lib/rocm.exp @@ -99,6 +99,56 @@ gdb_caching_proc allow_hipcc_tests {} { return 1 } +# ROCM_PATH is used by hipcc as well. +if {[info exists env(ROCM_PATH)]} { + set rocm_path $env(ROCM_PATH) +} else { + set rocm_path "/opt/rocm" +} + +# Get the gpu target to be passed to e.g., -mcpu=. +# +# If HCC_AMDGPU_TARGET is set in the environment, use it. Otherwise, +# try reading it from the system using the rocm_agent_enumerator +# utility. + +proc hcc_amdgpu_target {} { + if {![info exists ::gdb_hip_gpu]} { + # Look for HCC_AMDGPU_TARGET (same env var hipcc uses). If + # that fails, try using rocm_agent_enumerator (again, same as + # hipcc does). + if {[info exists env(HCC_AMDGPU_TARGET)]} { + set targets = $env(HCC_AMDGPU_TARGET); + } else { + set result \ + [remote_exec host \ + "${::rocm_path}/bin/rocm_agent_enumerator -t GPU"] + if {[lindex $result 0] != 0} { + error "rocm_agent_enumerator failed" + } + set targets [lindex $result 1] + } + + set ::gdb_hip_gpu "" + foreach val $targets { + # Ignore the 'gfx000' target reported by + # rocm_agent_enumerator. + if {$val != "gfx000"} { + set ::gdb_hip_gpu $val + break + } + } + + if {$::gdb_hip_gpu == ""} { + error "No valid AMD GPU target specified or found.\ + Please specify a valid target using the\ + HCC_AMDGPU_TARGET environment variable." + } + } + + return $::gdb_hip_gpu +} + # The lock file used to ensure that only one GDB has access to the GPU # at a time. set gpu_lock_filename $objdir/gpu-parallel.lock @@ -186,3 +236,12 @@ proc hip_devices_support_debug_multi_process {} { } return 1 } + +# Return true if the device supports precise memory. + +proc hip_device_supports_precise_memory {} { + set target [hcc_amdgpu_target] + set unsupported_targets \ + {gfx900 gfx906 gfx908 gfx1010 gfx1011 gfx1012 gfx1030 gfx1031 gfx1032} + return [expr [lsearch -exact $unsupported_targets $target] == -1] +} -- 2.42.0 ^ permalink raw reply [flat|nested] 10+ messages in thread
* Re: [PATCH 2/2] gdb/amdgpu: add precise-memory support 2023-09-13 15:28 ` [PATCH 2/2] gdb/amdgpu: add precise-memory support Simon Marchi @ 2023-09-13 21:32 ` Lancelot SIX 2023-09-14 15:51 ` Simon Marchi 2023-09-14 14:10 ` Tom Tromey 1 sibling, 1 reply; 10+ messages in thread From: Lancelot SIX @ 2023-09-13 21:32 UTC (permalink / raw) To: Simon Marchi; +Cc: gdb-patches Hi Simon, Thanks for doing this. I have a couple of remarks inlined in the patch. Best, Lancelot. On Wed, Sep 13, 2023 at 11:28:38AM -0400, Simon Marchi via Gdb-patches wrote: > The amd-dbgapi library exposes a setting called "memory precision" for > AMD GPUs [1]. Here's a copy of the description of the setting: > > The AMD GPU can overlap the execution of memory instructions with other > instructions. This can result in a wave stopping due to a memory violation > or hardware data watchpoint hit with a program counter beyond the > instruction that caused the wave to stop. > > Some architectures allow the hardware to be configured to always wait for > memory operations to complete before continuing. This will result in the > wave stopping at the instruction immediately after the one that caused the > stop event. Enabling this mode can make execution of waves significantly > slower. > > Expose this option through a new "amdgpu precise-memory" setting. > > The precise memory setting is per inferior. The setting is transferred > from one inferior to another when using the clone-inferior command, or > when a new inferior is created following an exec or a fork. > > It can be set before starting the inferior, in which case GDB will > attempt to apply what the user wants when attaching amd-dbgapi. If the > user has requested to enable precise memory, but it can't be enabled > (not all hardware supports it), GDB prints a warning. > > If precise memory is disabled, GDB prints a warning when hitting a > memory exception (translated into GDB_SIGNAL_SEGV or GDB_SIGNAL_BUS), > saying that the stop location may not be precise. > > Note that the precise memory setting also affects memory watchpoint > reporting, but the watchpoint support for AMD GPUs hasn't been > upstreamed to GDB yet. When we do upstream watchpoint support, GDB will > produce a similar warning message when stopping due to a watchpoint if > precise memory is disabled. > > Add a handful of tests. Add a util proc > "hip_device_supports_precise_memory", which indicates if the device used > for testing supports that feature. To implement it, also add a new > "hcc_amdgpu_target" proc, to return the architecture of the device used > for testing. > > [1] https://github.com/ROCm-Developer-Tools/ROCdbgapi/blob/687374258a27b5aab1309a7e8ded719e2f1ed3b1/include/amd-dbgapi.h.in#L6300-L6317 > > Change-Id: Ife1a99c0e960513da375ced8f8afaf8e47a61b3f > --- > gdb/amd-dbgapi-target.c | 199 +++++++++++++++++- > gdb/doc/gdb.texinfo | 43 ++++ > gdb/testsuite/gdb.rocm/precise-memory-exec.c | 44 ++++ > .../gdb.rocm/precise-memory-exec.exp | 62 ++++++ > gdb/testsuite/gdb.rocm/precise-memory-fork.c | 41 ++++ > .../gdb.rocm/precise-memory-fork.exp | 54 +++++ > .../precise-memory-multi-inferiors.exp | 87 ++++++++ > .../precise-memory-warning-sigsegv.cpp | 33 +++ > .../precise-memory-warning-sigsegv.exp | 49 +++++ > gdb/testsuite/gdb.rocm/precise-memory.cpp | 32 +++ > gdb/testsuite/gdb.rocm/precise-memory.exp | 57 +++++ > gdb/testsuite/lib/rocm.exp | 59 ++++++ > 12 files changed, 755 insertions(+), 5 deletions(-) > create mode 100644 gdb/testsuite/gdb.rocm/precise-memory-exec.c > create mode 100644 gdb/testsuite/gdb.rocm/precise-memory-exec.exp > create mode 100644 gdb/testsuite/gdb.rocm/precise-memory-fork.c > create mode 100644 gdb/testsuite/gdb.rocm/precise-memory-fork.exp > create mode 100644 gdb/testsuite/gdb.rocm/precise-memory-multi-inferiors.exp > create mode 100644 gdb/testsuite/gdb.rocm/precise-memory-warning-sigsegv.cpp > create mode 100644 gdb/testsuite/gdb.rocm/precise-memory-warning-sigsegv.exp > create mode 100644 gdb/testsuite/gdb.rocm/precise-memory.cpp > create mode 100644 gdb/testsuite/gdb.rocm/precise-memory.exp > > diff --git a/gdb/amd-dbgapi-target.c b/gdb/amd-dbgapi-target.c > index 22c269b7992c..cfb935df1163 100644 > --- a/gdb/amd-dbgapi-target.c > +++ b/gdb/amd-dbgapi-target.c > @@ -23,6 +23,7 @@ > #include "amdgpu-tdep.h" > #include "async-event.h" > #include "cli/cli-cmds.h" > +#include "cli/cli-decode.h" > #include "cli/cli-style.h" > #include "inf-loop.h" > #include "inferior.h" > @@ -139,6 +140,17 @@ struct amd_dbgapi_inferior_info > Initialized to true, since that's the default in amd-dbgapi too. */ > bool forward_progress_required = true; > > + struct > + { > + /* Whether precise memory reporting is requested. */ > + bool requested = false; > + > + /* Whether precise memory was requested and successfully enabled by > + dbgapi (it may not be available for the current hardware, for > + instance). */ > + bool enabled = false; > + } precise_memory; > + > std::unordered_map<decltype (amd_dbgapi_breakpoint_id_t::handle), > struct breakpoint *> > breakpoint_map; > @@ -1326,6 +1338,36 @@ amd_dbgapi_target::stopped_by_hw_breakpoint () > return false; > } > > +/* Set the process's memory access reporting precision. > + > + The precision can be ::AMD_DBGAPI_MEMORY_PRECISION_PRECISE (waves wait for > + memory instructions to complete before executing further instructions), or > + ::AMD_DBGAPI_MEMORY_PRECISION_NONE (memory instructions execute normally). > + > + Returns true if the precision is supported by the architecture of all agents > + in the process, or false if at least one agent does not support the > + requested precision. > + > + An error is thrown if setting the precision results in a status other than > + ::AMD_DBGAPI_STATUS_SUCCESS or ::AMD_DBGAPI_STATUS_ERROR_NOT_SUPPORTED. */ > + Would it be simpler if this helper function received a bool parameter instead of the amd_dbgapi_memory_precision_t one? This could avoid repeating this amd_dbgapi_memory_precision_t memory_precision = (info->precise_memory.requested ? AMD_DBGAPI_MEMORY_PRECISION_PRECISE : AMD_DBGAPI_MEMORY_PRECISION_NONE); before calling it. > +static bool > +set_process_memory_precision (amd_dbgapi_process_id_t process_id, > + amd_dbgapi_memory_precision_t precision) > +{ > + amd_dbgapi_status_t status > + = amd_dbgapi_set_memory_precision (process_id, precision); > + > + if (status == AMD_DBGAPI_STATUS_ERROR_NOT_SUPPORTED) > + return false; > + > + if (status != AMD_DBGAPI_STATUS_SUCCESS) > + error (_("amd_dbgapi_set_memory_precision failed (%s)"), > + get_status_string (status)); > + > + return true; > +} > + > /* Make the amd-dbgapi library attach to the process behind INF. > > Note that this is unrelated to the "attach" GDB concept / command. > @@ -1399,6 +1441,16 @@ attach_amd_dbgapi (inferior *inf) > amd_dbgapi_debug_printf ("process_id = %" PRIu64 ", notifier fd = %d", > info->process_id.handle, info->notifier); > > + amd_dbgapi_memory_precision_t memory_precision > + = (info->precise_memory.requested > + ? AMD_DBGAPI_MEMORY_PRECISION_PRECISE > + : AMD_DBGAPI_MEMORY_PRECISION_NONE); > + if (set_process_memory_precision (info->process_id, memory_precision)) > + info->precise_memory.enabled = info->precise_memory.requested; > + else > + warning > + (_("AMDGPU precise memory access reporting could not be enabled.")); > + > /* 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 > target. */ > @@ -1443,8 +1495,10 @@ detach_amd_dbgapi (inferior *inf) > for (auto &&value : info->breakpoint_map) > delete_breakpoint (value.second); > > - /* Reset the amd_dbgapi_inferior_info. */ > + /* Reset the amd_dbgapi_inferior_info, except for precise_memory_mode. */ > + bool precise_memory_requested = info->precise_memory.requested; > *info = amd_dbgapi_inferior_info (inf); > + info->precise_memory.requested = precise_memory_requested; > > maybe_reset_amd_dbgapi (); > } > @@ -1668,6 +1722,22 @@ amd_dbgapi_target_inferior_created (inferior *inf) > attach_amd_dbgapi (inf); > } > > +/* Callback called when an inferior is cloned. */ > + > +static void > +amd_dbgapi_target_inferior_cloned (inferior *original_inferior, > + inferior *new_inferior) > +{ > + auto *orig_info = get_amd_dbgapi_inferior_info (original_inferior); > + auto *new_info = get_amd_dbgapi_inferior_info (new_inferior); > + > + /* At this point, the process is not started. Therefore it is sufficient to > + copy the precise memory request, it will be applied when the process > + starts. */ > + gdb_assert (new_info->process_id == AMD_DBGAPI_PROCESS_NONE); > + new_info->precise_memory.requested = orig_info->precise_memory.requested; > +} > + > /* inferior_execd observer. */ > > static void > @@ -1677,6 +1747,13 @@ amd_dbgapi_inferior_execd (inferior *exec_inf, inferior *follow_inf) > attached to the old process image, so we need to detach and re-attach to > the new process image. */ > detach_amd_dbgapi (exec_inf); > + > + /* If using "follow-exec-mode new", carry over the precise-memory setting > + to the new inferior (otherwise, FOLLOW_INF and ORIG_INF point to the same > + inferior, so this is a no-op). */ > + get_amd_dbgapi_inferior_info (follow_inf)->precise_memory.requested > + = get_amd_dbgapi_inferior_info (exec_inf)->precise_memory.requested; > + > attach_amd_dbgapi (follow_inf); > } > > @@ -1686,11 +1763,22 @@ static void > amd_dbgapi_inferior_forked (inferior *parent_inf, inferior *child_inf, > target_waitkind fork_kind) > { > - if (child_inf != nullptr && fork_kind != TARGET_WAITKIND_VFORKED) > + if (child_inf != nullptr) > { > - scoped_restore_current_thread restore_thread; > - switch_to_thread (*child_inf->threads ().begin ()); > - attach_amd_dbgapi (child_inf); > + /* Copy precise-memory requested value from parent to child. */ > + amd_dbgapi_inferior_info *parent_info > + = get_amd_dbgapi_inferior_info (parent_inf); > + amd_dbgapi_inferior_info *child_info > + = get_amd_dbgapi_inferior_info (child_inf); > + child_info->precise_memory.requested > + = parent_info->precise_memory.requested; > + > + if (fork_kind != TARGET_WAITKIND_VFORKED) > + { > + scoped_restore_current_thread restore_thread; > + switch_to_thread (*child_inf->threads ().begin ()); > + attach_amd_dbgapi (child_inf); > + } > } > } > > @@ -1785,6 +1873,29 @@ amd_dbgapi_remove_breakpoint_callback > return AMD_DBGAPI_STATUS_SUCCESS; > } > > +/* signal_received observer. */ > + > +static void > +amd_dbgapi_target_signal_received (gdb_signal sig) > +{ > + amd_dbgapi_inferior_info *info > + = get_amd_dbgapi_inferior_info (current_inferior ()); > + > + if (info->process_id == AMD_DBGAPI_PROCESS_NONE) > + return; > + > + if (!ptid_is_gpu (inferior_thread ()->ptid)) > + return; > + > + if (sig != GDB_SIGNAL_SEGV && sig != GDB_SIGNAL_BUS) > + return; > + > + if (!info->precise_memory.enabled) > + gdb_printf ("\ I think there should be a _() surrounding the string. > +Warning: precise memory violation signal reporting is not enabled, reported\n\ > +location may not be accurate. See \"show amdgpu precise-memory\".\n"); > +} > + > /* Style for some kinds of messages. */ > > static cli_style_option fatal_error_style > @@ -1853,6 +1964,62 @@ amd_dbgapi_target::close () > delete_async_event_handler (&amd_dbgapi_async_event_handler); > } > > +/* Callback for "show amdgpu precise-memory". */ > + > +static void > +show_precise_memory_mode (struct ui_file *file, int from_tty, > + struct cmd_list_element *c, const char *value) > +{ > + amd_dbgapi_inferior_info *info > + = get_amd_dbgapi_inferior_info (current_inferior ()); > + > + gdb_printf (file, > + _("AMDGPU precise memory access reporting is %s " > + "(currently %s).\n"), > + info->precise_memory.requested ? "on" : "off", > + info->precise_memory.enabled ? "enabled" : "disabled"); > +} > + > +/* Callback for "set amdgpu precise-memory". */ > + > +static void > +set_precise_memory_mode (bool value) > +{ > + amd_dbgapi_inferior_info *info > + = get_amd_dbgapi_inferior_info (current_inferior ()); > + > + info->precise_memory.requested = value; > + > + if (info->process_id != AMD_DBGAPI_PROCESS_NONE) > + { > + amd_dbgapi_memory_precision_t memory_precision > + = (info->precise_memory.requested > + ? AMD_DBGAPI_MEMORY_PRECISION_PRECISE > + : AMD_DBGAPI_MEMORY_PRECISION_NONE); > + > + if (set_process_memory_precision (info->process_id, memory_precision)) > + info->precise_memory.enabled = info->precise_memory.requested; > + else > + warning > + (_("AMDGPU precise memory access reporting could not be enabled.")); > + } > +} > + > +/* Return whether precise-memory is requested for the current inferior. */ > + > +static bool > +get_precise_memory_mode () > +{ > + amd_dbgapi_inferior_info *info > + = get_amd_dbgapi_inferior_info (current_inferior ()); > + > + return info->precise_memory.requested; > +} > + > +/* List of set/show amdgpu commands. */ > +struct cmd_list_element *set_amdgpu_list; > +struct cmd_list_element *show_amdgpu_list; > + > /* List of set/show debug amd-dbgapi-lib commands. */ > struct cmd_list_element *set_debug_amd_dbgapi_lib_list; > struct cmd_list_element *show_debug_amd_dbgapi_lib_list; > @@ -1960,6 +2127,10 @@ _initialize_amd_dbgapi_target () > amd_dbgapi_set_log_level (get_debug_amd_dbgapi_lib_log_level ()); > > /* Install observers. */ > + gdb::observers::inferior_cloned.attach (amd_dbgapi_target_inferior_cloned, > + "amd-dbgapi"); > + gdb::observers::signal_received.attach (amd_dbgapi_target_signal_received, > + "amd-dbgapi"); > gdb::observers::inferior_created.attach > (amd_dbgapi_target_inferior_created, > amd_dbgapi_target_inferior_created_observer_token, "amd-dbgapi"); > @@ -1968,6 +2139,24 @@ _initialize_amd_dbgapi_target () > gdb::observers::inferior_exit.attach (amd_dbgapi_inferior_exited, "amd-dbgapi"); > gdb::observers::inferior_pre_detach.attach (amd_dbgapi_inferior_pre_detach, "amd-dbgapi"); > > + add_basic_prefix_cmd ("amdgpu", no_class, > + _("Generic command for setting amdgpu flags."), > + &set_amdgpu_list, 0, &setlist); > + > + add_show_prefix_cmd ("amdgpu", no_class, > + _("Generic command for showing amdgpu flags."), > + &show_amdgpu_list, 0, &showlist); > + > + add_setshow_boolean_cmd ("precise-memory", no_class, > + _("Set precise-memory mode."), > + _("Show precise-memory mode."), _("\ > +If on, precise memory reporting is enabled if/when the inferior is running.\n\ > +If off (default), precise memory reporting is disabled."), > + set_precise_memory_mode, > + get_precise_memory_mode, > + show_precise_memory_mode, > + &set_amdgpu_list, &show_amdgpu_list); > + > add_basic_prefix_cmd ("amd-dbgapi-lib", no_class, > _("Generic command for setting amd-dbgapi library " > "debugging flags."), > diff --git a/gdb/doc/gdb.texinfo b/gdb/doc/gdb.texinfo > index 9b7e06f31566..fa91d72695e7 100644 > --- a/gdb/doc/gdb.texinfo > +++ b/gdb/doc/gdb.texinfo > @@ -26794,6 +26794,49 @@ either not mapped or accessed with incompatible permissions. > If a single instruction raises more than one signal, they will be > reported one at a time each time the wavefront is continued. > > +@subsubsection @acronym{AMD GPU} Memory Violation Reporting > + > +A wavefront can report memory violation events. However, the program > +location at which they are reported may be after the machine instruction > +that caused them. This can result in the reported source statement > +being incorrect. The following commands can be used to control this > +behavior: > + > +@table @code > + > +@kindex set amdgpu precise-memory > +@cindex AMD GPU precise memory event reporting > +@item set amdgpu precise-memory @var{mode} > +Controls how @acronym{AMD GPU} devices detect memory violations, where > +@var{mode} can be: > + > +@table @code > + > +@item off > +The program location may not be immediately after the instruction that > +caused the memory violation. This is the default. > + > +@item on > +Requests that the program location will be immediately after the > +instruction that caused a memory violation. Enabling this mode may make > +the @acronym{AMD GPU} device execution significantly slower as it has to > +wait for each memory operation to complete before executing the next > +instruction. > + > +@end table > + > +The @code{set amdgpu precise-memory} parameter is per-inferior. When an ^ Isn't the parameter name just "amdgpu precise-memory"? > +inferior forks or execs, or the user uses the @code{clone-inferior} command, > +and an inferior is created as a result, the newly created inferior inherits > +the parameter value of the original inferior. > + > +@kindex show amdgpu precise-memory > +@cindex AMD GPU precise memory event reporting > +@item show amdgpu precise-memory > +Displays the currently requested AMD GPU precise memory setting. > + > +@end table > + > @subsubsection @acronym{AMD GPU} Logging > > The @samp{set debug amd-dbgapi} command can be used > diff --git a/gdb/testsuite/gdb.rocm/precise-memory-exec.c b/gdb/testsuite/gdb.rocm/precise-memory-exec.c > new file mode 100644 > index 000000000000..f0659a63fc5a > --- /dev/null > +++ b/gdb/testsuite/gdb.rocm/precise-memory-exec.c > @@ -0,0 +1,44 @@ > +/* Copyright 2021-2023 Free Software Foundation, Inc. > + > + This file is part of GDB. > + > + 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 <unistd.h> > +#include <stdlib.h> > +#include <stdio.h> > + > +static void > +second (void) > +{ > +} > + > +int > +main (int argc, char **argv) > +{ > + if (argc == 1) > + { > + /* First invocation */ Should the comment end with ". "? > + int ret = execl (argv[0], argv[0], "Hello", NULL); > + perror ("exec"); > + abort (); > + } > + else > + { > + /* Second invocation */ Here also. > + second (); > + } > + > + return 0; > +} > diff --git a/gdb/testsuite/gdb.rocm/precise-memory-exec.exp b/gdb/testsuite/gdb.rocm/precise-memory-exec.exp > new file mode 100644 > index 000000000000..26be6cf72146 > --- /dev/null > +++ b/gdb/testsuite/gdb.rocm/precise-memory-exec.exp > @@ -0,0 +1,62 @@ > +# Copyright 2021-2023 Free Software Foundation, Inc. > + > +# This file is part of GDB. > + > +# 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 the "set amdgpu precise-memory" setting is inherited by an inferior > +# created following an exec. > + > +load_lib rocm.exp > + > +require allow_hipcc_tests > + > +if { ![istarget "*-linux*"] } then { > + continue > +} Should this test be integrated in allow_hipcc_test? This would avoid having to repeat it in multiple testcases (and all testcases do not have such guard). Also, I'm not sure if there is any non-linux configuration which can satisfy allow_hipcc_tests, which would make this test redundant. > + > +standard_testfile .c > + > +if {[build_executable "failed to prepare $testfile" $testfile $srcfile {debug}]} { > + return > +} > + > +proc do_test { follow-exec-mode } { > + clean_restart $::binfile > + > + with_rocm_gpu_lock { > + if ![runto_main] { > + return > + } > + > + # Set precise-memory on the inferior before exec. > + gdb_test "show amdgpu precise-memory" " is off.*" \ > + "show amdgpu precise-memory before set" > + gdb_test "set amdgpu precise-memory on" > + gdb_test "show amdgpu precise-memory" " is on.*" \ > + "show amdgpu precise-memory after set" > + > + # Continue past exec. The precise-memory setting should > + # be on. > + gdb_test_no_output "set follow-exec-mode ${follow-exec-mode}" > + gdb_test "break second" > + gdb_test "continue" "Breakpoint 1(\.$::decimal)?, main .*" > + gdb_test "show amdgpu precise-memory" " is on.*" \ > + "show amdgpu precise-memory after exec" > + } > +} > + > +foreach_with_prefix follow-exec-mode {same new} { > + do_test ${follow-exec-mode} > +} > diff --git a/gdb/testsuite/gdb.rocm/precise-memory-fork.c b/gdb/testsuite/gdb.rocm/precise-memory-fork.c > new file mode 100644 > index 000000000000..67ce09f2c3dc > --- /dev/null > +++ b/gdb/testsuite/gdb.rocm/precise-memory-fork.c > @@ -0,0 +1,41 @@ > +/* Copyright 2021-2023 Free Software Foundation, Inc. > + > + This file is part of GDB. > + > + 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 <unistd.h> > + > +static void > +parent (void) > +{ > +} > + > +static void > +child (void) > +{ > +} > + > +int > +main (void) > +{ > + int pid = fork (); > + > + if (pid != 0) > + parent (); > + else > + child (); > + > + return 0; > +} > diff --git a/gdb/testsuite/gdb.rocm/precise-memory-fork.exp b/gdb/testsuite/gdb.rocm/precise-memory-fork.exp > new file mode 100644 > index 000000000000..0dc88b89f8a7 > --- /dev/null > +++ b/gdb/testsuite/gdb.rocm/precise-memory-fork.exp > @@ -0,0 +1,54 @@ > +# Copyright 2021-2023 Free Software Foundation, Inc. > + > +# This file is part of GDB. > + > +# 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 the "set amdgpu precise-memory" setting is inherited by a fork > +# child. > + > +load_lib rocm.exp > + > +require allow_hipcc_tests > + > +if { ![istarget "*-linux*"] } then { > + continue > +} Same remark here. > + > +standard_testfile .c > + > +if {[prepare_for_testing "failed to prepare $testfile" $testfile $srcfile {debug}]} { > + return > +} > + > +with_rocm_gpu_lock { > + if ![runto_main] { > + return > + } > + > + # Set precise-memory on in the parent, before fork. > + gdb_test "show amdgpu precise-memory" " is off.*" \ > + "show amdgpu precise-memory before set" > + gdb_test "set amdgpu precise-memory on" > + gdb_test "show amdgpu precise-memory" " is on.*" \ > + "show amdgpu precise-memory after set" > + > + # Continue past fork, following the child. The precise-memory setting should > + # be on. > + gdb_test "set follow-fork-mode child" > + gdb_test "break child" > + gdb_test "continue" "Thread 2.1 .* hit Breakpoint .*" > + gdb_test "show amdgpu precise-memory" " is on.*" \ > + "show amdgpu precise-memory after fork" > +} > diff --git a/gdb/testsuite/gdb.rocm/precise-memory-multi-inferiors.exp b/gdb/testsuite/gdb.rocm/precise-memory-multi-inferiors.exp > new file mode 100644 > index 000000000000..9968b422b0ee > --- /dev/null > +++ b/gdb/testsuite/gdb.rocm/precise-memory-multi-inferiors.exp > @@ -0,0 +1,87 @@ > +# Copyright 2021-2023 Free Software Foundation, Inc. > + > +# This file is part of GDB. > + > +# 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 the "set amdgpu precise-memory" setting is per-inferior, and > +# inherited by an inferior created using the clone-inferior command. > + > +load_lib rocm.exp > + > +require allow_hipcc_tests > + > +clean_restart > + > +set test_python [allow_python_tests] > + > +proc test_per_inferior { } { > + gdb_test "show amdgpu precise-memory" \ > + "AMDGPU precise memory access reporting is off \\(currently disabled\\)." \ > + "show initial value, inferior 1" > + if $::test_python { > + gdb_test "python print(gdb.parameter(\"amdgpu precise-memory\"))" \ > + "False" \ > + "show initial value using Python, inferior 1" > + } > + gdb_test_no_output "set amdgpu precise-memory" \ > + "set on inferior 1" > + gdb_test "show amdgpu precise-memory" \ > + "AMDGPU precise memory access reporting is on \\(currently disabled\\)." \ > + "show new value, inferior 1" > + if $::test_python { > + gdb_test "python print(gdb.parameter(\"amdgpu precise-memory\"))" \ > + "True" \ > + "show new value using Python, inferior 1" > + } > + > + gdb_test "add-inferior" "Added inferior 2" > + gdb_test "inferior 2" "Switching to inferior 2 .*" > + > + gdb_test "show amdgpu precise-memory" \ > + "AMDGPU precise memory access reporting is off \\(currently disabled\\)." \ > + "show initial value, inferior 2" > + if $::test_python { > + gdb_test "python print(gdb.parameter(\"amdgpu precise-memory\"))" \ > + "False" \ > + "show initial value using Python, inferior 2" > + } > +} > + > +proc test_copy_precise_memory_on_clone {precise_memory} { > + set value $precise_memory > + if {$precise_memory == "unspecified"} { > + set value off > + } > + > + clean_restart > + gdb_test "show amdgpu precise-memory" "is off.*" \ > + "show default amdgpu precise-memory" > + if {$precise_memory != "unspecified"} { > + gdb_test_no_output "set amdgpu precise-memory $value" > + gdb_test "show amdgpu precise-memory" "is $value.*" \ > + "show amdgpu precise-memory on original inferior" > + } > + > + gdb_test "clone-inferior" "Added inferior 2.*" > + gdb_test "inferior 2" > + gdb_test "show amdgpu precise-memory" "is $value.*" \ > + "show amdgpu precise-memory on cloned inferior" > +} > + > +test_per_inferior > + > +foreach_with_prefix precise_memory { unspecified on off } { > + test_copy_precise_memory_on_clone $precise_memory > +} > diff --git a/gdb/testsuite/gdb.rocm/precise-memory-warning-sigsegv.cpp b/gdb/testsuite/gdb.rocm/precise-memory-warning-sigsegv.cpp > new file mode 100644 > index 000000000000..58339e5391a6 > --- /dev/null > +++ b/gdb/testsuite/gdb.rocm/precise-memory-warning-sigsegv.cpp > @@ -0,0 +1,33 @@ > +/* Copyright 2021-2023 Free Software Foundation, Inc. > + > + This file is part of GDB. > + > + 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 <hip/hip_runtime.h> > + > +__global__ void > +kernel () > +{ > + int *p = nullptr; > + *p = 1; > +} > + > +int > +main (int argc, char* argv[]) > +{ > + hipLaunchKernelGGL (kernel, dim3 (1), dim3 (1), 0, 0); I think the "modern" way to write this would be: kernel<<<1, 1>>> (); This is mostly a remark, I don't mind using hipLaunchKernelGGL too much either. > + hipDeviceSynchronize (); > + return 0; > +} > diff --git a/gdb/testsuite/gdb.rocm/precise-memory-warning-sigsegv.exp b/gdb/testsuite/gdb.rocm/precise-memory-warning-sigsegv.exp > new file mode 100644 > index 000000000000..22e1f6eda254 > --- /dev/null > +++ b/gdb/testsuite/gdb.rocm/precise-memory-warning-sigsegv.exp > @@ -0,0 +1,49 @@ > +# Copyright 2021-2023 Free Software Foundation, Inc. > + > +# This file is part of GDB. > + > +# 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 when "amdgpu precise-memory" is off, hitting a SIGSEGV shows a > +# warning about the stop location maybe being inaccurate. > + > +load_lib rocm.exp > + > +require allow_hipcc_tests > + > +if { ![istarget "*-linux*"] } then { > + continue > +} Same remark here. > + > +standard_testfile .cpp > + > +if {[build_executable "failed to prepare" $testfile $srcfile {debug hip}]} { > + return > +} > + > +proc do_test { } { > + clean_restart $::binfile > + > + with_rocm_gpu_lock { > + if ![runto_main] { > + return > + } > + > + gdb_test_no_output "set amdgpu precise-memory off" > + gdb_test "continue" \ > + "SIGSEGV, Segmentation fault.*Warning: precise memory violation signal reporting is not enabled.*" > + } > +} > + > +do_test > diff --git a/gdb/testsuite/gdb.rocm/precise-memory.cpp b/gdb/testsuite/gdb.rocm/precise-memory.cpp > new file mode 100644 > index 000000000000..6e0a4d9bc63b > --- /dev/null > +++ b/gdb/testsuite/gdb.rocm/precise-memory.cpp > @@ -0,0 +1,32 @@ > +/* Copyright 2021-2023 Free Software Foundation, Inc. > + > + This file is part of GDB. > + > + 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 <hip/hip_runtime.h> > + > +__global__ void > +kernel () > +{ > + __builtin_amdgcn_s_sleep (1); > +} > + > +int > +main (int argc, char* argv[]) > +{ > + hipLaunchKernelGGL (kernel, dim3 (1), dim3 (1), 0, 0); Same, maybe prefer the kernel<<<1, 1>>> notation. > + hipDeviceSynchronize (); > + return 0; > +} > diff --git a/gdb/testsuite/gdb.rocm/precise-memory.exp b/gdb/testsuite/gdb.rocm/precise-memory.exp > new file mode 100644 > index 000000000000..bd2b12747c6f > --- /dev/null > +++ b/gdb/testsuite/gdb.rocm/precise-memory.exp > @@ -0,0 +1,57 @@ > +# Copyright 2022-2023 Free Software Foundation, Inc. > + > +# This file is part of GDB. > + > +# 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 showing the "amdgpu precise-memory" setting. > + > +load_lib rocm.exp > + > +require allow_hipcc_tests > + > +standard_testfile .cpp > + > +if {[build_executable "failed to prepare" $testfile $srcfile {debug hip}]} { > + return > +} > + > +proc do_test { } { > + clean_restart $::binfile > + > + with_rocm_gpu_lock { > + if ![runto_main] { > + return > + } > + > + gdb_test "show amdgpu precise-memory" \ > + "AMDGPU precise memory access reporting is off \\(currently disabled\\)." \ > + "show precise-memory setting in CLI before" > + > + if {[hip_device_supports_precise_memory]} { > + gdb_test_no_output "set amdgpu precise-memory on" > + set cli_effective_value "enabled" > + } else { > + gdb_test "set amdgpu precise-memory on" \ > + "warning: AMDGPU precise memory access reporting could not be enabled." > + set cli_effective_value "disabled" > + } > + > + gdb_test "show amdgpu precise-memory" \ > + "AMDGPU precise memory access reporting is on \\(currently ${cli_effective_value}\\)." \ > + "show precise-memory setting in CLI after" > + } > +} > + > +do_test > diff --git a/gdb/testsuite/lib/rocm.exp b/gdb/testsuite/lib/rocm.exp > index 98a3b308228d..22b294a5efae 100644 > --- a/gdb/testsuite/lib/rocm.exp > +++ b/gdb/testsuite/lib/rocm.exp > @@ -99,6 +99,56 @@ gdb_caching_proc allow_hipcc_tests {} { > return 1 > } > > +# ROCM_PATH is used by hipcc as well. > +if {[info exists env(ROCM_PATH)]} { > + set rocm_path $env(ROCM_PATH) > +} else { > + set rocm_path "/opt/rocm" > +} > + > +# Get the gpu target to be passed to e.g., -mcpu=. > +# > +# If HCC_AMDGPU_TARGET is set in the environment, use it. Otherwise, > +# try reading it from the system using the rocm_agent_enumerator > +# utility. > + > +proc hcc_amdgpu_target {} { There is a hcc_amdgpu_targets proc which enumerates the architecture of each agent present on the system. This is a fairly recent addition, it might have been introduced after you prepared this series. > + if {![info exists ::gdb_hip_gpu]} { > + # Look for HCC_AMDGPU_TARGET (same env var hipcc uses). If > + # that fails, try using rocm_agent_enumerator (again, same as > + # hipcc does). > + if {[info exists env(HCC_AMDGPU_TARGET)]} { > + set targets = $env(HCC_AMDGPU_TARGET); > + } else { > + set result \ > + [remote_exec host \ > + "${::rocm_path}/bin/rocm_agent_enumerator -t GPU"] > + if {[lindex $result 0] != 0} { > + error "rocm_agent_enumerator failed" > + } > + set targets [lindex $result 1] > + } > + > + set ::gdb_hip_gpu "" > + foreach val $targets { > + # Ignore the 'gfx000' target reported by > + # rocm_agent_enumerator. > + if {$val != "gfx000"} { > + set ::gdb_hip_gpu $val > + break > + } > + } > + > + if {$::gdb_hip_gpu == ""} { > + error "No valid AMD GPU target specified or found.\ > + Please specify a valid target using the\ > + HCC_AMDGPU_TARGET environment variable." > + } > + } > + > + return $::gdb_hip_gpu > +} > + > # The lock file used to ensure that only one GDB has access to the GPU > # at a time. > set gpu_lock_filename $objdir/gpu-parallel.lock > @@ -186,3 +236,12 @@ proc hip_devices_support_debug_multi_process {} { > } > return 1 > } > + > +# Return true if the device supports precise memory. Using hcc_amdgpu_targets, you could have a function which checks that all agents in the system support precise memory, not just the first one detected. This will reflect how `set amdgpu precise-memory` works. > + > +proc hip_device_supports_precise_memory {} { > + set target [hcc_amdgpu_target] > + set unsupported_targets \ > + {gfx900 gfx906 gfx908 gfx1010 gfx1011 gfx1012 gfx1030 gfx1031 gfx1032} > + return [expr [lsearch -exact $unsupported_targets $target] == -1] > +} > -- > 2.42.0 > ^ permalink raw reply [flat|nested] 10+ messages in thread
* Re: [PATCH 2/2] gdb/amdgpu: add precise-memory support 2023-09-13 21:32 ` Lancelot SIX @ 2023-09-14 15:51 ` Simon Marchi 0 siblings, 0 replies; 10+ messages in thread From: Simon Marchi @ 2023-09-14 15:51 UTC (permalink / raw) To: Lancelot SIX; +Cc: gdb-patches >> @@ -1326,6 +1338,36 @@ amd_dbgapi_target::stopped_by_hw_breakpoint () >> return false; >> } >> >> +/* Set the process's memory access reporting precision. >> + >> + The precision can be ::AMD_DBGAPI_MEMORY_PRECISION_PRECISE (waves wait for >> + memory instructions to complete before executing further instructions), or >> + ::AMD_DBGAPI_MEMORY_PRECISION_NONE (memory instructions execute normally). >> + >> + Returns true if the precision is supported by the architecture of all agents >> + in the process, or false if at least one agent does not support the >> + requested precision. >> + >> + An error is thrown if setting the precision results in a status other than >> + ::AMD_DBGAPI_STATUS_SUCCESS or ::AMD_DBGAPI_STATUS_ERROR_NOT_SUPPORTED. */ >> + > > Would it be simpler if this helper function received a bool parameter > instead of the amd_dbgapi_memory_precision_t one? This could avoid > repeating this > > amd_dbgapi_memory_precision_t memory_precision > = (info->precise_memory.requested > ? AMD_DBGAPI_MEMORY_PRECISION_PRECISE > : AMD_DBGAPI_MEMORY_PRECISION_NONE); > > before calling it. Indeed. In fact, we can factor out more... I came up with: static void try_set_process_memory_precision (amd_dbgapi_inferior_info &info) { auto mode = (info.precise_memory.requested ? AMD_DBGAPI_MEMORY_PRECISION_PRECISE : AMD_DBGAPI_MEMORY_PRECISION_NONE); amd_dbgapi_status_t status = amd_dbgapi_set_memory_precision (info.process_id, mode); if (status == AMD_DBGAPI_STATUS_ERROR_NOT_SUPPORTED) warning (_("AMDGPU precise memory access reporting could not be enabled.")); else if (status != AMD_DBGAPI_STATUS_SUCCESS) error (_("amd_dbgapi_set_memory_precision failed (%s)"), get_status_string (status)); } ... such that callers just need to do: try_set_process_memory_precision (*info); I'll put that in v2, unless you think it's a bad idea. >> @@ -1785,6 +1873,29 @@ amd_dbgapi_remove_breakpoint_callback >> return AMD_DBGAPI_STATUS_SUCCESS; >> } >> >> +/* signal_received observer. */ >> + >> +static void >> +amd_dbgapi_target_signal_received (gdb_signal sig) >> +{ >> + amd_dbgapi_inferior_info *info >> + = get_amd_dbgapi_inferior_info (current_inferior ()); >> + >> + if (info->process_id == AMD_DBGAPI_PROCESS_NONE) >> + return; >> + >> + if (!ptid_is_gpu (inferior_thread ()->ptid)) >> + return; >> + >> + if (sig != GDB_SIGNAL_SEGV && sig != GDB_SIGNAL_BUS) >> + return; >> + >> + if (!info->precise_memory.enabled) >> + gdb_printf ("\ > > I think there should be a _() surrounding the string. Done. >> +The @code{set amdgpu precise-memory} parameter is per-inferior. When an > ^ > Isn't the parameter name just "amdgpu precise-memory"? Yes, done. >> diff --git a/gdb/testsuite/gdb.rocm/precise-memory-exec.c b/gdb/testsuite/gdb.rocm/precise-memory-exec.c >> new file mode 100644 >> index 000000000000..f0659a63fc5a >> --- /dev/null >> +++ b/gdb/testsuite/gdb.rocm/precise-memory-exec.c >> @@ -0,0 +1,44 @@ >> +/* Copyright 2021-2023 Free Software Foundation, Inc. >> + >> + This file is part of GDB. >> + >> + 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 <unistd.h> >> +#include <stdlib.h> >> +#include <stdio.h> >> + >> +static void >> +second (void) >> +{ >> +} >> + >> +int >> +main (int argc, char **argv) >> +{ >> + if (argc == 1) >> + { >> + /* First invocation */ > > Should the comment end with ". "? Done. > >> + int ret = execl (argv[0], argv[0], "Hello", NULL); >> + perror ("exec"); >> + abort (); >> + } >> + else >> + { >> + /* Second invocation */ > > Here also. Done. > >> + second (); >> + } >> + >> + return 0; >> +} >> diff --git a/gdb/testsuite/gdb.rocm/precise-memory-exec.exp b/gdb/testsuite/gdb.rocm/precise-memory-exec.exp >> new file mode 100644 >> index 000000000000..26be6cf72146 >> --- /dev/null >> +++ b/gdb/testsuite/gdb.rocm/precise-memory-exec.exp >> @@ -0,0 +1,62 @@ >> +# Copyright 2021-2023 Free Software Foundation, Inc. >> + >> +# This file is part of GDB. >> + >> +# 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 the "set amdgpu precise-memory" setting is inherited by an inferior >> +# created following an exec. >> + >> +load_lib rocm.exp >> + >> +require allow_hipcc_tests >> + >> +if { ![istarget "*-linux*"] } then { >> + continue >> +} > > Should this test be integrated in allow_hipcc_test? This would avoid > having to repeat it in multiple testcases (and all testcases do not have > such guard). > > Also, I'm not sure if there is any non-linux configuration which can > satisfy allow_hipcc_tests, which would make this test redundant. Yeah I think we can put it in allow_hipcc_tests. If/when we want to run this test on, let's say, Windows, we will be able to change the check in allow_hipcc_tests. The existing tests in the upstream repo don't have the linux check, I guess there should have been some checks. I'll add a new preparatory patch in v2. >> diff --git a/gdb/testsuite/gdb.rocm/precise-memory-warning-sigsegv.cpp b/gdb/testsuite/gdb.rocm/precise-memory-warning-sigsegv.cpp >> new file mode 100644 >> index 000000000000..58339e5391a6 >> --- /dev/null >> +++ b/gdb/testsuite/gdb.rocm/precise-memory-warning-sigsegv.cpp >> @@ -0,0 +1,33 @@ >> +/* Copyright 2021-2023 Free Software Foundation, Inc. >> + >> + This file is part of GDB. >> + >> + 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 <hip/hip_runtime.h> >> + >> +__global__ void >> +kernel () >> +{ >> + int *p = nullptr; >> + *p = 1; >> +} >> + >> +int >> +main (int argc, char* argv[]) >> +{ >> + hipLaunchKernelGGL (kernel, dim3 (1), dim3 (1), 0, 0); > > I think the "modern" way to write this would be: > > kernel<<<1, 1>>> (); > > This is mostly a remark, I don't mind using hipLaunchKernelGGL too much > either. I think we wanted to migrate all to the new form, let's go with the new form. >> diff --git a/gdb/testsuite/lib/rocm.exp b/gdb/testsuite/lib/rocm.exp >> index 98a3b308228d..22b294a5efae 100644 >> --- a/gdb/testsuite/lib/rocm.exp >> +++ b/gdb/testsuite/lib/rocm.exp >> @@ -99,6 +99,56 @@ gdb_caching_proc allow_hipcc_tests {} { >> return 1 >> } >> >> +# ROCM_PATH is used by hipcc as well. >> +if {[info exists env(ROCM_PATH)]} { >> + set rocm_path $env(ROCM_PATH) >> +} else { >> + set rocm_path "/opt/rocm" >> +} >> + >> +# Get the gpu target to be passed to e.g., -mcpu=. >> +# >> +# If HCC_AMDGPU_TARGET is set in the environment, use it. Otherwise, >> +# try reading it from the system using the rocm_agent_enumerator >> +# utility. >> + >> +proc hcc_amdgpu_target {} { > > There is a hcc_amdgpu_targets proc which enumerates the architecture of > each agent present on the system. This is a fairly recent addition, it > might have been introduced after you prepared this series. Oh, that's true, no need to introduce hcc_amdgpu_target. >> @@ -186,3 +236,12 @@ proc hip_devices_support_debug_multi_process {} { >> } >> return 1 >> } >> + >> +# Return true if the device supports precise memory. > > Using hcc_amdgpu_targets, you could have a function which checks that > all agents in the system support precise memory, not just the first one > detected. This will reflect how `set amdgpu precise-memory` works. Oh, right, we have this downstream. I sync'ed the patch with the downstream code. Thanks, will send a v2. Simon ^ permalink raw reply [flat|nested] 10+ messages in thread
* Re: [PATCH 2/2] gdb/amdgpu: add precise-memory support 2023-09-13 15:28 ` [PATCH 2/2] gdb/amdgpu: add precise-memory support Simon Marchi 2023-09-13 21:32 ` Lancelot SIX @ 2023-09-14 14:10 ` Tom Tromey 2023-09-14 16:00 ` Simon Marchi 1 sibling, 1 reply; 10+ messages in thread From: Tom Tromey @ 2023-09-14 14:10 UTC (permalink / raw) To: Simon Marchi via Gdb-patches; +Cc: Simon Marchi >>>>> "Simon" == Simon Marchi via Gdb-patches <gdb-patches@sourceware.org> writes: Simon> It can be set before starting the inferior, in which case GDB will Simon> attempt to apply what the user wants when attaching amd-dbgapi. If the Simon> user has requested to enable precise memory, but it can't be enabled Simon> (not all hardware supports it), GDB prints a warning. I was curious if this setting is preserved by fork or exec and so ought to be copied when those events occur. Tom ^ permalink raw reply [flat|nested] 10+ messages in thread
* Re: [PATCH 2/2] gdb/amdgpu: add precise-memory support 2023-09-14 14:10 ` Tom Tromey @ 2023-09-14 16:00 ` Simon Marchi 2023-09-14 16:18 ` Tom Tromey 0 siblings, 1 reply; 10+ messages in thread From: Simon Marchi @ 2023-09-14 16:00 UTC (permalink / raw) To: Tom Tromey, Simon Marchi via Gdb-patches On 9/14/23 10:10, Tom Tromey wrote: >>>>>> "Simon" == Simon Marchi via Gdb-patches <gdb-patches@sourceware.org> writes: > > Simon> It can be set before starting the inferior, in which case GDB will > Simon> attempt to apply what the user wants when attaching amd-dbgapi. If the > Simon> user has requested to enable precise memory, but it can't be enabled > Simon> (not all hardware supports it), GDB prints a warning. > > I was curious if this setting is preserved by fork or exec and so ought > to be copied when those events occur. Yes, it is mentioned in the commit message, and there are tests for it in this patch. Or maybe I don't understand what you mean. Simon ^ permalink raw reply [flat|nested] 10+ messages in thread
* Re: [PATCH 2/2] gdb/amdgpu: add precise-memory support 2023-09-14 16:00 ` Simon Marchi @ 2023-09-14 16:18 ` Tom Tromey 2023-09-14 16:18 ` Simon Marchi 0 siblings, 1 reply; 10+ messages in thread From: Tom Tromey @ 2023-09-14 16:18 UTC (permalink / raw) To: Simon Marchi via Gdb-patches; +Cc: Tom Tromey, Simon Marchi >>>>> "Simon" == Simon Marchi via Gdb-patches <gdb-patches@sourceware.org> writes: >> I was curious if this setting is preserved by fork or exec and so ought >> to be copied when those events occur. Simon> Yes, it is mentioned in the commit message, and there are tests for it Simon> in this patch. Or maybe I don't understand what you mean. I just totally missed it...?! Sorry about that. Tom ^ permalink raw reply [flat|nested] 10+ messages in thread
* Re: [PATCH 2/2] gdb/amdgpu: add precise-memory support 2023-09-14 16:18 ` Tom Tromey @ 2023-09-14 16:18 ` Simon Marchi 0 siblings, 0 replies; 10+ messages in thread From: Simon Marchi @ 2023-09-14 16:18 UTC (permalink / raw) To: Tom Tromey, Simon Marchi via Gdb-patches On 9/14/23 12:18, Tom Tromey wrote: >>>>>> "Simon" == Simon Marchi via Gdb-patches <gdb-patches@sourceware.org> writes: > >>> I was curious if this setting is preserved by fork or exec and so ought >>> to be copied when those events occur. > > Simon> Yes, it is mentioned in the commit message, and there are tests for it > Simon> in this patch. Or maybe I don't understand what you mean. > > I just totally missed it...?! Sorry about that. > > Tom No problem, better be safe than sorry. Simon ^ permalink raw reply [flat|nested] 10+ messages in thread
* Re: [PATCH 1/2] gdb: add inferior_cloned observable 2023-09-13 15:28 [PATCH 1/2] gdb: add inferior_cloned observable Simon Marchi 2023-09-13 15:28 ` [PATCH 2/2] gdb/amdgpu: add precise-memory support Simon Marchi @ 2023-09-14 14:08 ` Tom Tromey 2023-09-14 15:53 ` Simon Marchi 1 sibling, 1 reply; 10+ messages in thread From: Tom Tromey @ 2023-09-14 14:08 UTC (permalink / raw) To: Simon Marchi via Gdb-patches; +Cc: Simon Marchi >>>>> "Simon" == Simon Marchi via Gdb-patches <gdb-patches@sourceware.org> writes: Simon> The following patch makes the amdgpu port transfer a property from the Simon> original inferior to the new inferior when using the clone-inferior Simon> command. Add the inferior_cloned observable to help with this. FWIW this patch seems totally fine to me. Tom ^ permalink raw reply [flat|nested] 10+ messages in thread
* Re: [PATCH 1/2] gdb: add inferior_cloned observable 2023-09-14 14:08 ` [PATCH 1/2] gdb: add inferior_cloned observable Tom Tromey @ 2023-09-14 15:53 ` Simon Marchi 0 siblings, 0 replies; 10+ messages in thread From: Simon Marchi @ 2023-09-14 15:53 UTC (permalink / raw) To: Tom Tromey, Simon Marchi via Gdb-patches On 9/14/23 10:08, Tom Tromey wrote: >>>>>> "Simon" == Simon Marchi via Gdb-patches <gdb-patches@sourceware.org> writes: > > Simon> The following patch makes the amdgpu port transfer a property from the > Simon> original inferior to the new inferior when using the clone-inferior > Simon> command. Add the inferior_cloned observable to help with this. > > FWIW this patch seems totally fine to me. > > Tom Thanks, I'll add you Approved-By. Simon ^ permalink raw reply [flat|nested] 10+ messages in thread
end of thread, other threads:[~2023-09-14 16:18 UTC | newest] Thread overview: 10+ messages (download: mbox.gz / follow: Atom feed) -- links below jump to the message on this page -- 2023-09-13 15:28 [PATCH 1/2] gdb: add inferior_cloned observable Simon Marchi 2023-09-13 15:28 ` [PATCH 2/2] gdb/amdgpu: add precise-memory support Simon Marchi 2023-09-13 21:32 ` Lancelot SIX 2023-09-14 15:51 ` Simon Marchi 2023-09-14 14:10 ` Tom Tromey 2023-09-14 16:00 ` Simon Marchi 2023-09-14 16:18 ` Tom Tromey 2023-09-14 16:18 ` Simon Marchi 2023-09-14 14:08 ` [PATCH 1/2] gdb: add inferior_cloned observable Tom Tromey 2023-09-14 15:53 ` Simon Marchi
This is a public inbox, see mirroring instructions for how to clone and mirror all data and code used for this inbox; as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).