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

[Simon Marchi <simon.marchi@efficios.com>]

>> @@ -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.
Indeed.  In fact, we can factor out more... I came up with:

    static void
    try_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_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));
    }

... such that callers just need to do:

    try_set_process_memory_precision (*info);

I'll put that in v2, unless you think it's a bad idea.

>> @@ -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.

Done.

>> +The @code{set amdgpu precise-memory} parameter is per-inferior.  When an
>              ^
> Isn't the parameter name just "amdgpu precise-memory"?

Yes, done.

>> 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 ".  "?

Done.

> 
>> +      int ret = execl (argv[0], argv[0], "Hello", NULL);
>> +      perror ("exec");
>> +      abort ();
>> +    }
>> +  else
>> +    {
>> +      /* Second invocation */
> 
> Here also.

Done.

> 
>> +      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.

Yeah I think we can put it in allow_hipcc_tests.  If/when we want to run
this test on, let's say, Windows, we will be able to change the check in
allow_hipcc_tests.

The existing tests in the upstream repo don't have the linux check, I
guess there should have been some checks.  I'll add a new preparatory
patch in v2.

>> 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.

I think we wanted to migrate all to the new form, let's go with the new
form.

>> 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.

Oh, that's true, no need to introduce hcc_amdgpu_target.

>> @@ -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.

Oh, right, we have this downstream.  I sync'ed the patch with the
downstream code.

Thanks, will send a v2.

Simon