diff options
Diffstat (limited to 'gdb')
| -rw-r--r-- | gdb/amd-dbgapi-target.c | 182 | ||||
| -rw-r--r-- | gdb/doc/gdb.texinfo | 43 | ||||
| -rw-r--r-- | gdb/testsuite/gdb.rocm/precise-memory-exec.c | 44 | ||||
| -rw-r--r-- | gdb/testsuite/gdb.rocm/precise-memory-exec.exp | 58 | ||||
| -rw-r--r-- | gdb/testsuite/gdb.rocm/precise-memory-fork.c | 41 | ||||
| -rw-r--r-- | gdb/testsuite/gdb.rocm/precise-memory-fork.exp | 50 | ||||
| -rw-r--r-- | gdb/testsuite/gdb.rocm/precise-memory-multi-inferiors.exp | 87 | ||||
| -rw-r--r-- | gdb/testsuite/gdb.rocm/precise-memory-warning-sigsegv.cpp | 33 | ||||
| -rw-r--r-- | gdb/testsuite/gdb.rocm/precise-memory-warning-sigsegv.exp | 45 | ||||
| -rw-r--r-- | gdb/testsuite/gdb.rocm/precise-memory.cpp | 32 | ||||
| -rw-r--r-- | gdb/testsuite/gdb.rocm/precise-memory.exp | 57 | ||||
| -rw-r--r-- | gdb/testsuite/lib/rocm.exp | 19 |
12 files changed, 683 insertions, 8 deletions
diff --git a/gdb/amd-dbgapi-target.c b/gdb/amd-dbgapi-target.c index 8eafcea..e578358 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" @@ -116,9 +117,12 @@ get_amd_dbgapi_target_inferior_created_observer_token () struct amd_dbgapi_inferior_info { - explicit amd_dbgapi_inferior_info (inferior *inf) + explicit amd_dbgapi_inferior_info (inferior *inf, + bool precise_memory_requested = false) : inf (inf) - {} + { + precise_memory.requested = precise_memory_requested; + } /* Backlink to inferior. */ inferior *inf; @@ -139,6 +143,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; + + /* 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 +1341,31 @@ amd_dbgapi_target::stopped_by_hw_breakpoint () return false; } +/* Set the process' memory access reporting precision mode. + + Warn if the requested mode is not supported on at least one agent in the + process. + + Error out if setting the requested mode failed for some other reason. */ + +static void +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_SUCCESS) + info.precise_memory.enabled = info.precise_memory.requested; + else 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)); +} + /* Make the amd-dbgapi library attach to the process behind INF. Note that this is unrelated to the "attach" GDB concept / command. @@ -1399,6 +1439,8 @@ attach_amd_dbgapi (inferior *inf) amd_dbgapi_debug_printf ("process_id = %" PRIu64 ", notifier fd = %d", info->process_id.handle, info->notifier); + set_process_memory_precision (*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 target. */ @@ -1443,8 +1485,8 @@ detach_amd_dbgapi (inferior *inf) for (auto &&value : info->breakpoint_map) delete_breakpoint (value.second); - /* Reset the amd_dbgapi_inferior_info. */ - *info = amd_dbgapi_inferior_info (inf); + /* Reset the amd_dbgapi_inferior_info, except for precise_memory_mode. */ + *info = amd_dbgapi_inferior_info (inf, info->precise_memory.requested); maybe_reset_amd_dbgapi (); } @@ -1668,6 +1710,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 +1735,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 +1751,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 +1861,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 +1952,51 @@ 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) + set_process_memory_precision (*info); +} + +/* 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 +2104,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 +2116,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 9b7e06f..aa3c677 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{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 0000000..a1f941d --- /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 0000000..601f30b --- /dev/null +++ b/gdb/testsuite/gdb.rocm/precise-memory-exec.exp @@ -0,0 +1,58 @@ +# 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 + +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 0000000..67ce09f --- /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 0000000..347b62b --- /dev/null +++ b/gdb/testsuite/gdb.rocm/precise-memory-fork.exp @@ -0,0 +1,50 @@ +# 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 + +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 0000000..9968b42 --- /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 0000000..bf1451a --- /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[]) +{ + kernel<<<1, 1>>> (); + 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 0000000..2813ca7 --- /dev/null +++ b/gdb/testsuite/gdb.rocm/precise-memory-warning-sigsegv.exp @@ -0,0 +1,45 @@ +# 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 + +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 0000000..034f023 --- /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[]) +{ + kernel<<<1, 1>>> (); + 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 0000000..62b7515 --- /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_devices_support_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 791f1b4..fcdf665 100644 --- a/gdb/testsuite/lib/rocm.exp +++ b/gdb/testsuite/lib/rocm.exp @@ -190,3 +190,22 @@ proc hip_devices_support_debug_multi_process {} { } return 1 } + +# Return true if all the devices on the host support precise memory. + +proc hip_devices_support_precise_memory {} { + set unsupported_targets \ + {gfx900 gfx906 gfx908 gfx1010 gfx1011 gfx1012 gfx1030 gfx1031 gfx1032} + + set targets [hcc_amdgpu_targets] + if { [llength $targets] == 0 } { + return 0 + } + + foreach target $targets { + if { [lsearch -exact $unsupported_targets $target] != -1 } { + return 0 + } + } + return 1 +} |