From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from lndn.lancelotsix.com (lndn.lancelotsix.com [51.195.220.111]) by sourceware.org (Postfix) with ESMTPS id F27EE3858C5F for ; Wed, 13 Sep 2023 21:32:56 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org F27EE3858C5F Authentication-Results: sourceware.org; dmarc=pass (p=reject dis=none) header.from=lancelotsix.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=lancelotsix.com Received: from octopus (cust120-dsl54.idnet.net [212.69.54.120]) by lndn.lancelotsix.com (Postfix) with ESMTPSA id E413280910; Wed, 13 Sep 2023 21:32:55 +0000 (UTC) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/simple; d=lancelotsix.com; s=2021; t=1694640776; bh=30JrLI/mcjZfftH5JOTnaAu5OMvDzUIP6HCSW+4fALk=; h=Date:From:To:Cc:Subject:References:In-Reply-To:From; b=HnsJnH2tmIF03AjkXAAtBE970e5CI3t1yeG1KPFkBeA1szV63nBDmUuu4IlBIc5lq KA3nq1qcq0Iaq6E1z8l/ujjwWcEoyQBwmJ8wrk0189ujWQ/+XCbuLHwS51TrfSdI1V owpUCZ6s0F342U1hJ9k1sW30zWU7Dp+JUtAHYHKw76rd67kRk8s7wkwfNfRMPNdwn+ uquMatYuIWVAIo2lq5sDO/BYtTwr36MOyIIOwqFDOtrcexCek1GHtQ/29lz43swlc3 okipLf/ub9wTPfQwW7azW3kECdyTt8rM+CG4i6rhEDpKXghCztMJFuVpdV9IVlw6DF QPbwspIJ+NGvA== Date: Wed, 13 Sep 2023 22:32:50 +0100 From: Lancelot SIX To: Simon Marchi Cc: gdb-patches@sourceware.org Subject: Re: [PATCH 2/2] gdb/amdgpu: add precise-memory support Message-ID: <20230913213250.cwgwrowfrawelfac@octopus> References: <20230913152845.1540064-1-simon.marchi@efficios.com> <20230913152845.1540064-2-simon.marchi@efficios.com> MIME-Version: 1.0 Content-Type: text/plain; charset=us-ascii Content-Disposition: inline In-Reply-To: <20230913152845.1540064-2-simon.marchi@efficios.com> X-Greylist: Sender succeeded SMTP AUTH, not delayed by milter-greylist-4.6.2 (lndn.lancelotsix.com [0.0.0.0]); Wed, 13 Sep 2023 21:32:56 +0000 (UTC) X-Spam-Status: No, score=-11.8 required=5.0 tests=BAYES_00,DKIM_SIGNED,DKIM_VALID,DKIM_VALID_AU,DKIM_VALID_EF,GIT_PATCH_0,KAM_SHORT,SPF_HELO_NONE,SPF_PASS,TXREP autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org List-Id: 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 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 . */ > + > +#include > +#include > +#include > + > +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 . > + > +# 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 . */ > + > +#include > + > +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 . > + > +# 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 . > + > +# 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 . */ > + > +#include > + > +__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 . > + > +# 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 . */ > + > +#include > + > +__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 . > + > +# 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 >