@@ -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,8 +117,9 @@ get_amd_dbgapi_target_inferior_created_observer_token ()
struct amd_dbgapi_inferior_info
{
- explicit amd_dbgapi_inferior_info (inferior *inf)
- : inf (inf)
+ explicit amd_dbgapi_inferior_info (inferior *inf,
+ bool precise_memory_requested = false)
+ : inf {inf}, precise_memory {precise_memory_requested}
{}
/* Backlink to inferior. */
@@ -139,6 +141,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 +1339,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 +1437,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 +1483,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 +1708,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 +1733,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 +1749,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 +1859,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 +1950,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 +2102,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 +2114,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."),
@@ -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
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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}
+}
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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"
+}
new file mode 100644
@@ -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
+}
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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
@@ -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
+}