diff --git a/gdb/amd-dbgapi-target.c b/gdb/amd-dbgapi-target.c
index 22c269b7992c..507824decf8d 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,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."),
diff --git a/gdb/doc/gdb.texinfo b/gdb/doc/gdb.texinfo
index 9b7e06f31566..aa3c67788871 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 000000000000..a1f941dae030
--- /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..601f30b3be8a
--- /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 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..347b62ba88a6
--- /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 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..bf1451ad8d5b
--- /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 000000000000..2813ca781f2a
--- /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 000000000000..034f02382adc
--- /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 000000000000..62b7515c251b
--- /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 791f1b403bb7..fcdf665aef96 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
+}
