Re: [PATCH 2/2] gdb/amdgpu: add precise-memory support

[Lancelot SIX <lsix@lancelotsix.com>]

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
>