FreeChainXenon/binutils-gdb
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions

gdb: initial support for ROCm platform (AMDGPU) debugging

Browse source 
This patch adds the foundation for GDB to be able to debug programs
offloaded to AMD GPUs using the AMD ROCm platform [1].  The latest
public release of the ROCm release at the time of writing is 5.4, so
this is what this patch targets.

The ROCm platform allows host programs to schedule bits of code for
execution on GPUs or similar accelerators.  The programs running on GPUs
are typically referred to as `kernels` (not related to operating system
kernels).

Programs offloaded with the AMD ROCm platform can be written in the HIP
language [2], OpenCL and OpenMP, but we're going to focus on HIP here.
The HIP language consists of a C++ Runtime API and kernel language.
Here's an example of a very simple HIP program:

    #include "hip/hip_runtime.h"
    #include <cassert>

    __global__ void
    do_an_addition (int a, int b, int *out)
    {
      *out = a + b;
    }

    int
    main ()
    {
      int *result_ptr, result;

      /* Allocate memory for the device to write the result to.  */
      hipError_t error = hipMalloc (&result_ptr, sizeof (int));
      assert (error == hipSuccess);

      /* Run `do_an_addition` on one workgroup containing one work item.  */
      do_an_addition<<<dim3(1), dim3(1), 0, 0>>> (1, 2, result_ptr);

      /* Copy result from device to host.  Note that this acts as a synchronization
         point, waiting for the kernel dispatch to complete.  */
      error = hipMemcpyDtoH (&result, result_ptr, sizeof (int));
      assert (error == hipSuccess);

      printf ("result is %d\n", result);
      assert (result == 3);

      return 0;
    }

This program can be compiled with:

    $ hipcc simple.cpp -g -O0 -o simple

... where `hipcc` is the HIP compiler, shipped with ROCm releases.  This
generates an ELF binary for the host architecture, containing another
ELF binary with the device code.  The ELF for the device can be
inspected with:

    $ roc-obj-ls simple
    1       host-x86_64-unknown-linux                                           file://simple#offset=8192&size=0
    1       hipv4-amdgcn-amd-amdhsa--gfx906                                     file://simple#offset=8192&size=34216
    $ roc-obj-extract 'file://simple#offset=8192&size=34216'
    $ file simple-offset8192-size34216.co
    simple-offset8192-size34216.co: ELF 64-bit LSB shared object, *unknown arch 0xe0* version 1, dynamically linked, with debug_info, not stripped
                                                                                 ^
                       amcgcn architecture that my `file` doesn't know about ----´

Running the program gives the very unimpressive result:

    $ ./simple
    result is 3

While running, this host program has copied the device program into the
GPU's memory and spawned an execution thread on it.  The goal of this
GDB port is to let the user debug host threads and these GPU threads
simultaneously.  Here's a sample session using a GDB with this patch
applied:

    $ ./gdb -q -nx --data-directory=data-directory ./simple
    Reading symbols from ./simple...
    (gdb) break do_an_addition
    Function "do_an_addition" not defined.
    Make breakpoint pending on future shared library load? (y or [n]) y
    Breakpoint 1 (do_an_addition) pending.
    (gdb) r
    Starting program: /home/smarchi/build/binutils-gdb-amdgpu/gdb/simple
    [Thread debugging using libthread_db enabled]
    Using host libthread_db library "/lib/x86_64-linux-gnu/libthread_db.so.1".
    [New Thread 0x7ffff5db7640 (LWP 1082911)]
    [New Thread 0x7ffef53ff640 (LWP 1082913)]
    [Thread 0x7ffef53ff640 (LWP 1082913) exited]
    [New Thread 0x7ffdecb53640 (LWP 1083185)]
    [New Thread 0x7ffff54bf640 (LWP 1083186)]
    [Thread 0x7ffdecb53640 (LWP 1083185) exited]
    [Switching to AMDGPU Wave 2:2:1:1 (0,0,0)/0]

    Thread 6 hit Breakpoint 1, do_an_addition (a=<error reading variable: DWARF-2 expression error: `DW_OP_regx' operations must be used either alone or in conjunction with DW_OP_piece or DW_OP_bit_piece.>,
        b=<error reading variable: DWARF-2 expression error: `DW_OP_regx' operations must be used either alone or in conjunction with DW_OP_piece or DW_OP_bit_piece.>,
        out=<error reading variable: DWARF-2 expression error: `DW_OP_regx' operations must be used either alone or in conjunction with DW_OP_piece or DW_OP_bit_piece.>) at simple.cpp:24
    24        *out = a + b;
    (gdb) info inferiors
      Num  Description       Connection           Executable
    * 1    process 1082907   1 (native)           /home/smarchi/build/binutils-gdb-amdgpu/gdb/simple
    (gdb) info threads
      Id   Target Id                                    Frame
      1    Thread 0x7ffff5dc9240 (LWP 1082907) "simple" 0x00007ffff5e9410b in ?? () from /opt/rocm-5.4.0/lib/libhsa-runtime64.so.1
      2    Thread 0x7ffff5db7640 (LWP 1082911) "simple" __GI___ioctl (fd=3, request=3222817548) at ../sysdeps/unix/sysv/linux/ioctl.c:36
      5    Thread 0x7ffff54bf640 (LWP 1083186) "simple" __GI___ioctl (fd=3, request=3222817548) at ../sysdeps/unix/sysv/linux/ioctl.c:36
    * 6    AMDGPU Wave 2:2:1:1 (0,0,0)/0                do_an_addition (
        a=<error reading variable: DWARF-2 expression error: `DW_OP_regx' operations must be used either alone or in conjunction with DW_OP_piece or DW_OP_bit_piece.>,
        b=<error reading variable: DWARF-2 expression error: `DW_OP_regx' operations must be used either alone or in conjunction with DW_OP_piece or DW_OP_bit_piece.>,
        out=<error reading variable: DWARF-2 expression error: `DW_OP_regx' operations must be used either alone or in conjunction with DW_OP_piece or DW_OP_bit_piece.>) at simple.cpp:24
    (gdb) bt
    Python Exception <class 'gdb.error'>: Unhandled dwarf expression opcode 0xe1
    #0  do_an_addition (a=<error reading variable: DWARF-2 expression error: `DW_OP_regx' operations must be used either alone or in conjunction with DW_OP_piece or DW_OP_bit_piece.>,
        b=<error reading variable: DWARF-2 expression error: `DW_OP_regx' operations must be used either alone or in conjunction with DW_OP_piece or DW_OP_bit_piece.>,
        out=<error reading variable: DWARF-2 expression error: `DW_OP_regx' operations must be used either alone or in conjunction with DW_OP_piece or DW_OP_bit_piece.>) at simple.cpp:24
    (gdb) continue
    Continuing.
    result is 3
    warning: Temporarily disabling breakpoints for unloaded shared library "file:///home/smarchi/build/binutils-gdb-amdgpu/gdb/simple#offset=8192&size=67208"
    [Thread 0x7ffff54bf640 (LWP 1083186) exited]
    [Thread 0x7ffff5db7640 (LWP 1082911) exited]
    [Inferior 1 (process 1082907) exited normally]

One thing to notice is the host and GPU threads appearing under
the same inferior.  This is a design goal for us, as programmers tend to
think of the threads running on the GPU as part of the same program as
the host threads, so showing them in the same inferior in GDB seems
natural.  Also, the host and GPU threads share a global memory space,
which fits the inferior model.

Another thing to notice is the error messages when trying to read
variables or printing a backtrace.  This is expected for the moment,
since the AMD GPU compiler produces some DWARF that uses some
non-standard extensions:

  https://llvm.org/docs/AMDGPUDwarfExtensionsForHeterogeneousDebugging.html

There were already some patches posted by Zoran Zaric earlier to make
GDB support these extensions:

  https://inbox.sourceware.org/gdb-patches/20211105113849.118800-1-zoran.zaric@amd.com/

We think it's better to get the basic support for AMD GPU in first,
which will then give a better justification for GDB to support these
extensions.

GPU threads are named `AMDGPU Wave`: a wave is essentially a hardware
thread using the SIMT (single-instruction, multiple-threads) [3]
execution model.

GDB uses the amd-dbgapi library [4], included in the ROCm platform, for
a few things related to AMD GPU threads debugging.  Different components
talk to the library, as show on the following diagram:

    +---------------------------+     +-------------+     +------------------+
    | GDB   | amd-dbgapi target | <-> |     AMD     |     |    Linux kernel  |
    |       +-------------------+     |   Debugger  |     +--------+         |
    |       | amdgcn gdbarch    | <-> |     API     | <=> | AMDGPU |         |
    |       +-------------------+     |             |     | driver |         |
    |       | solib-rocm        | <-> | (dbgapi.so) |     +--------+---------+
    +---------------------------+     +-------------+

  - The amd-dbgapi target is a target_ops implementation used to control
    execution of GPU threads.  While the debugging of host threads works
    by using the ptrace / wait Linux kernel interface (as usual), control
    of GPU threads is done through a special interface (dubbed `kfd`)
    exposed by the `amdgpu` Linux kernel module.  GDB doesn't interact
    directly with `kfd`, but instead goes through the amd-dbgapi library
    (AMD Debugger API on the diagram).

    Since it provides execution control, the amd-dbgapi target should
    normally be a process_stratum_target, not just a target_ops.  More
    on that later.

  - The amdgcn gdbarch (describing the hardware architecture of the GPU
    execution units) offloads some requests to the amd-dbgapi library,
    so that knowledge about the various architectures doesn't need to be
    duplicated and baked in GDB.  This is for example for things like
    the list of registers.

  - The solib-rocm component is an solib provider that fetches the list of
    code objects loaded on the device from the amd-dbgapi library, and
    makes GDB read their symbols.  This is very similar to other solib
    providers that handle shared libraries, except that here the shared
    libraries are the pieces of code loaded on the device.

Given that Linux host threads are managed by the linux-nat target, and
the GPU threads are managed by the amd-dbgapi target, having all threads
appear in the same inferior requires the two targets to be in that
inferior's target stack.  However, there can only be one
process_stratum_target in a given target stack, since there can be only
one target per slot.  To achieve it, we therefore resort the hack^W
solution of placing the amd-dbgapi target in the arch_stratum slot of
the target stack, on top of the linux-nat target.  Doing so allows the
amd-dbgapi target to intercept target calls and handle them if they
concern GPU threads, and offload to beneath otherwise.  See
amd_dbgapi_target::fetch_registers for a simple example:

    void
    amd_dbgapi_target::fetch_registers (struct regcache *regcache, int regno)
    {
      if (!ptid_is_gpu (regcache->ptid ()))
        {
          beneath ()->fetch_registers (regcache, regno);
          return;
        }

      // handle it
    }

ptids of GPU threads are crafted with the following pattern:

  (pid, 1, wave id)

Where pid is the inferior's pid and "wave id" is the wave handle handed
to us by the amd-dbgapi library (in practice, a monotonically
incrementing integer).  The idea is that on Linux systems, the
combination (pid != 1, lwp == 1) is not possible.  lwp == 1 would always
belong to the init process, which would also have pid == 1 (and it's
improbable for the init process to offload work to the GPU and much less
for the user to debug it).  We can therefore differentiate GPU and
non-GPU ptids this way.  See ptid_is_gpu for more details.

Note that we believe that this scheme could break down in the context of
containers, where the initial process executed in a container has pid 1
(in its own pid namespace).  For instance, if you were to execute a ROCm
program in a container, then spawn a GDB in that container and attach to
the process, it will likely not work.  This is a known limitation.  A
workaround for this is to have a dummy process (like a shell) fork and
execute the program of interest.

The amd-dbgapi target watches native inferiors, and "attaches" to them
using amd_dbgapi_process_attach, which gives it a notifier fd that is
registered in the event loop (see enable_amd_dbgapi).  Note that this
isn't the same "attach" as in PTRACE_ATTACH, but being ptrace-attached
is a precondition for amd_dbgapi_process_attach to work.  When the
debugged process enables the ROCm runtime, the amd-dbgapi target gets
notified through that fd, and pushes itself on the target stack of the
inferior.  The amd-dbgapi target is then able to intercept target_ops
calls.  If the debugged process disables the ROCm runtime, the
amd-dbgapi target unpushes itself from the target stack.

This way, the amd-dbgapi target's footprint stays minimal when debugging
a process that doesn't use the AMD ROCm platform, it does not intercept
target calls.

The amd-dbgapi library is found using pkg-config.  Since enabling
support for the amdgpu architecture (amdgpu-tdep.c) depends on the
amd-dbgapi library being present, we have the following logic for
the interaction with --target and --enable-targets:

 - if the user explicitly asks for amdgcn support with
   --target=amdgcn-*-* or --enable-targets=amdgcn-*-*, we probe for
   the amd-dbgapi and fail if not found

 - if the user uses --enable-targets=all, we probe for amd-dbgapi,
   enable amdgcn support if found, disable amdgcn support if not found

 - if the user uses --enable-targets=all and --with-amd-dbgapi=yes,
   we probe for amd-dbgapi, enable amdgcn if found and fail if not found

 - if the user uses --enable-targets=all and --with-amd-dbgapi=no,
   we do not probe for amd-dbgapi, disable amdgcn support

 - otherwise, amd-dbgapi is not probed for and support for amdgcn is not
   enabled

Finally, a simple test is included.  It only tests hitting a breakpoint
in device code and resuming execution, pretty much like the example
shown above.

[1] https://docs.amd.com/category/ROCm_v5.4
[2] https://docs.amd.com/bundle/HIP-Programming-Guide-v5.4
[3] https://en.wikipedia.org/wiki/Single_instruction,_multiple_threads
[4] https://docs.amd.com/bundle/ROCDebugger-API-Guide-v5.4

Change-Id: I591edca98b8927b1e49e4b0abe4e304765fed9ee
Co-Authored-By: Zoran Zaric <zoran.zaric@amd.com>
Co-Authored-By: Laurent Morichetti <laurent.morichetti@amd.com>
Co-Authored-By: Tony Tye <Tony.Tye@amd.com>
Co-Authored-By: Lancelot SIX <lancelot.six@amd.com>
Co-Authored-By: Pedro Alves <pedro@palves.net>
This commit is contained in:
Simon Marchi 2023-01-03 15:07:07 -05:00
parent 8b588f4212
commit 18b4d0736b
18 changed files with 5155 additions and 138 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

17
gdb/Makefile.in
View file

@ -227,6 +227,9 @@ PTHREAD_LIBS = @PTHREAD_LIBS@
DEBUGINFOD_CFLAGS = @DEBUGINFOD_CFLAGS@
DEBUGINFOD_LIBS = @DEBUGINFOD_LIBS@
AMD_DBGAPI_CFLAGS = @AMD_DBGAPI_CFLAGS@
AMD_DBGAPI_LIBS = @AMD_DBGAPI_LIBS@
RDYNAMIC = @RDYNAMIC@
# Where is the INTL library? Typically in ../intl.
@ -633,7 +636,8 @@ INTERNAL_CFLAGS_BASE = \
$(ZSTD_CFLAGS) $(BFD_CFLAGS) $(INCLUDE_CFLAGS) $(LIBDECNUMBER_CFLAGS) \
$(INTL_CFLAGS) $(INCGNU) $(INCSUPPORT) $(LIBBACKTRACE_INC) \
$(ENABLE_CFLAGS) $(INTERNAL_CPPFLAGS) $(SRCHIGH_CFLAGS) \
$(TOP_CFLAGS) $(PTHREAD_CFLAGS) $(DEBUGINFOD_CFLAGS) $(GMPINC)
$(TOP_CFLAGS) $(PTHREAD_CFLAGS) $(DEBUGINFOD_CFLAGS) $(GMPINC) \
$(AMD_DBGAPI_CFLAGS)
INTERNAL_WARN_CFLAGS = $(INTERNAL_CFLAGS_BASE) $(GDB_WARN_CFLAGS)
INTERNAL_CFLAGS = $(INTERNAL_WARN_CFLAGS) $(GDB_WERROR_CFLAGS)
@ -655,7 +659,7 @@ INTERNAL_LDFLAGS = \
CLIBS = $(SIM) $(READLINE) $(OPCODES) $(LIBCTF) $(BFD) $(ZLIB) $(ZSTD_LIBS) \
$(LIBSUPPORT) $(INTL) $(LIBIBERTY) $(LIBDECNUMBER) \
$(XM_CLIBS) $(GDBTKLIBS) $(LIBBACKTRACE_LIB) \
@LIBS@ @GUILE_LIBS@ @PYTHON_LIBS@ \
@LIBS@ @GUILE_LIBS@ @PYTHON_LIBS@ $(AMD_DBGAPI_LIBS) \
$(LIBEXPAT) $(LIBLZMA) $(LIBBABELTRACE) $(LIBIPT) \
$(WIN32LIBS) $(LIBGNU) $(LIBGNU_EXTRA_LIBS) $(LIBICONV) \
$(GMPLIBS) $(SRCHIGH_LIBS) $(LIBXXHASH) $(PTHREAD_LIBS) \
@ -693,6 +697,12 @@ SIM_OBS = @SIM_OBS@
# Target-dependent object files.
TARGET_OBS = @TARGET_OBS@
# All target-dependent object files that require the amd-dbgapi
# target to be available (used with --enable-targets=all).
ALL_AMD_DBGAPI_TARGET_OBS = \
amdgpu-tdep.o \
solib-rocm.o
# All target-dependent objects files that require 64-bit CORE_ADDR
# (used with --enable-targets=all --enable-64-bit-bfd).
ALL_64_TARGET_OBS = \
@ -1637,6 +1647,7 @@ ALLDEPFILES = \
alpha-netbsd-tdep.c \
alpha-obsd-tdep.c \
alpha-tdep.c \
amd-dbgapi-target.c \
amd64-bsd-nat.c \
amd64-darwin-tdep.c \
amd64-dicos-tdep.c \
@ -1652,6 +1663,7 @@ ALLDEPFILES = \
amd64-ravenscar-thread.c \
amd64-sol2-tdep.c \
amd64-tdep.c \
amdgpu-tdep.c \
arc-linux-nat.c \
arc-tdep.c \
arm-bsd-tdep.c \
@ -1793,6 +1805,7 @@ ALLDEPFILES = \
sh-tdep.c \
sol2-tdep.c \
solib-aix.c \
solib-rocm.c \
solib-svr4.c \
sparc-linux-nat.c \
sparc-linux-tdep.c \

7
gdb/NEWS
View file

@ -244,6 +244,8 @@ GNU/Linux/LoongArch (gdbserver) loongarch*-*-linux*
GNU/Linux/CSKY (gdbserver) csky*-*linux*
AMDGPU amdgcn-*-*
* MI changes
** The async record stating the stopped reason 'breakpoint-hit' now
@ -338,6 +340,11 @@ GNU/Linux/CSKY (gdbserver) csky*-*linux*
GDB now supports floating-point on LoongArch GNU/Linux.
* AMD GPU ROCm debugging support
GDB now supports debugging programs offloaded to AMD GPUs using the ROCm
platform.
*** Changes in GDB 12
* DBX mode is deprecated, and will be removed in GDB 13

15
gdb/README
View file

@ -541,6 +541,21 @@ more obscure GDB `configure' options are not listed here.
speeds up various GDB operations such as symbol loading. Enabled
by default if libxxhash is found.
`--with-amd-dbgapi=[auto,yes,no]'
Whether to use the amd-dbgapi library to support local debugging of
AMD GCN architecture GPUs.
When explicitly requesting support for an AMD GCN architecture through
`--enable-targets' or `--target', there is no need to use
`--with-amd-dbgapi': `configure' will automatically look for the
amd-dbgapi library and fail if not found.
When using --enable-targets=all, support for the AMD GCN architecture will
only be included if the amd-dbgapi is found. `--with-amd-dbgapi=yes' can
be used to make it a failure if the amd-dbgapi library is not found.
`--with-amd-dbgapi=no' can be used to prevent looking for the amd-dbgapi
library altogether.
`--without-included-regex'
Don't use the regex library included with GDB (as part of the
libiberty library). This is the default on hosts with version 2

1966
gdb/amd-dbgapi-target.c Normal file
View file

File diff suppressed because it is too large Load diff

116
gdb/amd-dbgapi-target.h Normal file
View file

@ -0,0 +1,116 @@
/* Target used to communicate with the AMD Debugger API.
Copyright (C) 2019-2022 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/>. */
#ifndef AMD_DBGAPI_TARGET_H
#define AMD_DBGAPI_TARGET_H 1
#include "gdbsupport/observable.h"
#include <amd-dbgapi/amd-dbgapi.h>
struct inferior;
namespace detail
{
template <typename T>
using is_amd_dbgapi_handle
= gdb::Or<std::is_same<T, amd_dbgapi_address_class_id_t>,
std::is_same<T, amd_dbgapi_address_space_id_t>,
std::is_same<T, amd_dbgapi_architecture_id_t>,
std::is_same<T, amd_dbgapi_agent_id_t>,
std::is_same<T, amd_dbgapi_breakpoint_id_t>,
std::is_same<T, amd_dbgapi_code_object_id_t>,
std::is_same<T, amd_dbgapi_dispatch_id_t>,
std::is_same<T, amd_dbgapi_displaced_stepping_id_t>,
std::is_same<T, amd_dbgapi_event_id_t>,
std::is_same<T, amd_dbgapi_process_id_t>,
std::is_same<T, amd_dbgapi_queue_id_t>,
std::is_same<T, amd_dbgapi_register_class_id_t>,
std::is_same<T, amd_dbgapi_register_id_t>,
std::is_same<T, amd_dbgapi_watchpoint_id_t>,
std::is_same<T, amd_dbgapi_wave_id_t>>;
} /* namespace detail */
/* Get the token of amd-dbgapi's inferior_created observer. */
const gdb::observers::token &
get_amd_dbgapi_target_inferior_created_observer_token ();
/* Comparison operators for amd-dbgapi handle types. */
template <typename T,
typename = gdb::Requires<detail::is_amd_dbgapi_handle<T>>>
bool
operator== (const T &lhs, const T &rhs)
{
return lhs.handle == rhs.handle;
}
template <typename T,
typename = gdb::Requires<detail::is_amd_dbgapi_handle<T>>>
bool
operator!= (const T &lhs, const T &rhs)
{
return !(lhs == rhs);
}
/* Return true if the given ptid is a GPU thread (wave) ptid. */
static inline bool
ptid_is_gpu (ptid_t ptid)
{
/* FIXME: Currently using values that are known not to conflict with other
processes to indicate if it is a GPU thread. ptid.pid 1 is the init
process and is the only process that could have a ptid.lwp of 1. The init
process cannot have a GPU. No other process can have a ptid.lwp of 1.
The GPU wave ID is stored in the ptid.tid. */
return ptid.pid () != 1 && ptid.lwp () == 1;
}
/* Return INF's amd_dbgapi process id. */
amd_dbgapi_process_id_t get_amd_dbgapi_process_id (inferior *inf);
/* Get the amd-dbgapi wave id for PTID. */
static inline amd_dbgapi_wave_id_t
get_amd_dbgapi_wave_id (ptid_t ptid)
{
gdb_assert (ptid_is_gpu (ptid));
return amd_dbgapi_wave_id_t {
static_cast<decltype (amd_dbgapi_wave_id_t::handle)> (ptid.tid ())
};
}
/* Get the textual version of STATUS.
Always returns non-nullptr, and asserts that STATUS has a valid value. */
static inline const char *
get_status_string (amd_dbgapi_status_t status)
{
const char *ret;
status = amd_dbgapi_get_status_string (status, &ret);
gdb_assert (status == AMD_DBGAPI_STATUS_SUCCESS);
return ret;
}
#endif /* AMD_DBGAPI_TARGET_H */

1367
gdb/amdgpu-tdep.c Normal file
View file

File diff suppressed because it is too large Load diff

93
gdb/amdgpu-tdep.h Normal file
View file

@ -0,0 +1,93 @@
/* Target-dependent code for the AMDGPU architectures.
Copyright (C) 2019-2022 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/>. */
#ifndef AMDGPU_TDEP_H
#define AMDGPU_TDEP_H
#include "gdbarch.h"
#include <amd-dbgapi/amd-dbgapi.h>
#include <unordered_map>
/* Provide std::unordered_map::Hash for amd_dbgapi_register_id_t. */
struct register_id_hash
{
size_t
operator() (const amd_dbgapi_register_id_t &register_id) const
{
return std::hash<decltype (register_id.handle)> () (register_id.handle);
}
};
/* Provide std::unordered_map::Equal for amd_dbgapi_register_id_t. */
struct register_id_equal_to
{
bool
operator() (const amd_dbgapi_register_id_t &lhs,
const amd_dbgapi_register_id_t &rhs) const
{
return std::equal_to<decltype (lhs.handle)> () (lhs.handle, rhs.handle);
}
};
/* AMDGPU architecture specific information. */
struct amdgpu_gdbarch_tdep : gdbarch_tdep_base
{
/* This architecture's breakpoint instruction. */
gdb::unique_xmalloc_ptr<gdb_byte> breakpoint_instruction_bytes;
size_t breakpoint_instruction_size;
/* A vector of register_ids indexed by their equivalent gdb regnum. */
std::vector<amd_dbgapi_register_id_t> register_ids;
/* A vector of register_properties indexed by their equivalent gdb regnum. */
std::vector<amd_dbgapi_register_properties_t> register_properties;
/* A vector of register names indexed by their equivalent gdb regnum. */
std::vector<std::string> register_names;
/* A vector of register types created from the amd-dbgapi type strings,
indexed by their equivalent gdb regnum. These are computed lazily by
amdgpu_register_type, entries that haven't been computed yet are
nullptr. */
std::vector<type *> register_types;
/* A vector of GDB register numbers indexed by DWARF register number.
Unused DWARF register numbers map to value -1. */
std::vector<int> dwarf_regnum_to_gdb_regnum;
/* A map of gdb regnums keyed by they equivalent register_id. */
std::unordered_map<amd_dbgapi_register_id_t, int, register_id_hash,
register_id_equal_to>
regnum_map;
/* A map of register_class_ids keyed by their name. */
std::unordered_map<std::string, amd_dbgapi_register_class_id_t>
register_class_map;
};
/* Return true if GDBARCH is of an AMDGPU architecture. */
bool is_amdgpu_arch (struct gdbarch *gdbarch);
/* Return the amdgpu-specific data associated to ARCH. */
amdgpu_gdbarch_tdep *get_amdgpu_gdbarch_tdep (gdbarch *arch);
#endif /* AMDGPU_TDEP_H */

425
gdb/configure vendored
View file

@ -770,11 +770,10 @@ PKGVERSION
CODESIGN_CERT
DEBUGINFOD_LIBS
DEBUGINFOD_CFLAGS
PKG_CONFIG_LIBDIR
PKG_CONFIG_PATH
PKG_CONFIG
HAVE_NATIVE_GCORE_TARGET
TARGET_OBS
AMD_DBGAPI_LIBS
AMD_DBGAPI_CFLAGS
ENABLE_BFD_64_BIT_FALSE
ENABLE_BFD_64_BIT_TRUE
subdirs
@ -796,6 +795,9 @@ INCINTL
LIBINTL_DEP
LIBINTL
USE_NLS
PKG_CONFIG_LIBDIR
PKG_CONFIG_PATH
PKG_CONFIG
CCDEPMODE
DEPDIR
am__leading_dot
@ -909,6 +911,7 @@ with_auto_load_dir
with_auto_load_safe_path
enable_targets
enable_64_bit_bfd
with_amd_dbgapi
enable_gdbmi
enable_tui
enable_gdbtk
@ -975,11 +978,13 @@ CXXFLAGS
CCC
CPP
CXXCPP
MAKEINFO
MAKEINFOFLAGS
PKG_CONFIG
PKG_CONFIG_PATH
PKG_CONFIG_LIBDIR
MAKEINFO
MAKEINFOFLAGS
AMD_DBGAPI_CFLAGS
AMD_DBGAPI_LIBS
DEBUGINFOD_CFLAGS
DEBUGINFOD_LIBS
YACC
@ -1668,6 +1673,7 @@ Optional Packages:
[--with-auto-load-dir]
--without-auto-load-safe-path
do not restrict auto-loaded files locations
--with-amd-dbgapi support for the amd-dbgapi target (yes / no / auto)
--with-debuginfod Enable debuginfo lookups with debuginfod
(auto/yes/no)
--with-libunwind-ia64 use libunwind frame unwinding for ia64 targets
@ -1734,14 +1740,18 @@ Some influential environment variables:
CXXFLAGS C++ compiler flags
CPP C preprocessor
CXXCPP C++ preprocessor
MAKEINFO Parent configure detects if it is of sufficient version.
MAKEINFOFLAGS
Parameters for MAKEINFO.
PKG_CONFIG path to pkg-config utility
PKG_CONFIG_PATH
directories to add to pkg-config's search path
PKG_CONFIG_LIBDIR
path overriding pkg-config's built-in search path
MAKEINFO Parent configure detects if it is of sufficient version.
MAKEINFOFLAGS
Parameters for MAKEINFO.
AMD_DBGAPI_CFLAGS
C compiler flags for AMD_DBGAPI, overriding pkg-config
AMD_DBGAPI_LIBS
linker flags for AMD_DBGAPI, overriding pkg-config
DEBUGINFOD_CFLAGS
C compiler flags for DEBUGINFOD, overriding pkg-config
DEBUGINFOD_LIBS
@ -11439,7 +11449,7 @@ else
lt_dlunknown=0; lt_dlno_uscore=1; lt_dlneed_uscore=2
lt_status=$lt_dlunknown
cat > conftest.$ac_ext <<_LT_EOF
#line 11442 "configure"
#line 11452 "configure"
#include "confdefs.h"
#if HAVE_DLFCN_H
@ -11545,7 +11555,7 @@ else
lt_dlunknown=0; lt_dlno_uscore=1; lt_dlneed_uscore=2
lt_status=$lt_dlunknown
cat > conftest.$ac_ext <<_LT_EOF
#line 11548 "configure"
#line 11558 "configure"
#include "confdefs.h"
#if HAVE_DLFCN_H
@ -17525,6 +17535,130 @@ else CCDEPMODE=depmode=$am_cv_CC_dependencies_compiler_type
fi
# Since the first call to PKG_CHECK_MODULES may not happen (is guarded by
# a condition), we must call PKG_PROG_PKG_CONFIG explicitly to probe for
# pkg-config.
if test "x$ac_cv_env_PKG_CONFIG_set" != "xset"; then
if test -n "$ac_tool_prefix"; then
# Extract the first word of "${ac_tool_prefix}pkg-config", so it can be a program name with args.
set dummy ${ac_tool_prefix}pkg-config; ac_word=$2
{ $as_echo "$as_me:${as_lineno-$LINENO}: checking for $ac_word" >&5
$as_echo_n "checking for $ac_word... " >&6; }
if ${ac_cv_path_PKG_CONFIG+:} false; then :
$as_echo_n "(cached) " >&6
else
case $PKG_CONFIG in
[\\/]* | ?:[\\/]*)
ac_cv_path_PKG_CONFIG="$PKG_CONFIG" # Let the user override the test with a path.
;;
*)
as_save_IFS=$IFS; IFS=$PATH_SEPARATOR
for as_dir in $PATH
do
IFS=$as_save_IFS
test -z "$as_dir" && as_dir=.
for ac_exec_ext in '' $ac_executable_extensions; do
if as_fn_executable_p "$as_dir/$ac_word$ac_exec_ext"; then
ac_cv_path_PKG_CONFIG="$as_dir/$ac_word$ac_exec_ext"
$as_echo "$as_me:${as_lineno-$LINENO}: found $as_dir/$ac_word$ac_exec_ext" >&5
break 2
fi
done
done
IFS=$as_save_IFS
;;
esac
fi
PKG_CONFIG=$ac_cv_path_PKG_CONFIG
if test -n "$PKG_CONFIG"; then
{ $as_echo "$as_me:${as_lineno-$LINENO}: result: $PKG_CONFIG" >&5
$as_echo "$PKG_CONFIG" >&6; }
else
{ $as_echo "$as_me:${as_lineno-$LINENO}: result: no" >&5
$as_echo "no" >&6; }
fi
fi
if test -z "$ac_cv_path_PKG_CONFIG"; then
ac_pt_PKG_CONFIG=$PKG_CONFIG
# Extract the first word of "pkg-config", so it can be a program name with args.
set dummy pkg-config; ac_word=$2
{ $as_echo "$as_me:${as_lineno-$LINENO}: checking for $ac_word" >&5
$as_echo_n "checking for $ac_word... " >&6; }
if ${ac_cv_path_ac_pt_PKG_CONFIG+:} false; then :
$as_echo_n "(cached) " >&6
else
case $ac_pt_PKG_CONFIG in
[\\/]* | ?:[\\/]*)
ac_cv_path_ac_pt_PKG_CONFIG="$ac_pt_PKG_CONFIG" # Let the user override the test with a path.
;;
*)
as_save_IFS=$IFS; IFS=$PATH_SEPARATOR
for as_dir in $PATH
do
IFS=$as_save_IFS
test -z "$as_dir" && as_dir=.
for ac_exec_ext in '' $ac_executable_extensions; do
if as_fn_executable_p "$as_dir/$ac_word$ac_exec_ext"; then
ac_cv_path_ac_pt_PKG_CONFIG="$as_dir/$ac_word$ac_exec_ext"
$as_echo "$as_me:${as_lineno-$LINENO}: found $as_dir/$ac_word$ac_exec_ext" >&5
break 2
fi
done
done
IFS=$as_save_IFS
;;
esac
fi
ac_pt_PKG_CONFIG=$ac_cv_path_ac_pt_PKG_CONFIG
if test -n "$ac_pt_PKG_CONFIG"; then
{ $as_echo "$as_me:${as_lineno-$LINENO}: result: $ac_pt_PKG_CONFIG" >&5
$as_echo "$ac_pt_PKG_CONFIG" >&6; }
else
{ $as_echo "$as_me:${as_lineno-$LINENO}: result: no" >&5
$as_echo "no" >&6; }
fi
if test "x$ac_pt_PKG_CONFIG" = x; then
PKG_CONFIG=""
else
case $cross_compiling:$ac_tool_warned in
yes:)
{ $as_echo "$as_me:${as_lineno-$LINENO}: WARNING: using cross tools not prefixed with host triplet" >&5
$as_echo "$as_me: WARNING: using cross tools not prefixed with host triplet" >&2;}
ac_tool_warned=yes ;;
esac
PKG_CONFIG=$ac_pt_PKG_CONFIG
fi
else
PKG_CONFIG="$ac_cv_path_PKG_CONFIG"
fi
fi
if test -n "$PKG_CONFIG"; then
_pkg_min_version=0.9.0
{ $as_echo "$as_me:${as_lineno-$LINENO}: checking pkg-config is at least version $_pkg_min_version" >&5
$as_echo_n "checking pkg-config is at least version $_pkg_min_version... " >&6; }
if $PKG_CONFIG --atleast-pkgconfig-version $_pkg_min_version; then
{ $as_echo "$as_me:${as_lineno-$LINENO}: result: yes" >&5
$as_echo "yes" >&6; }
else
{ $as_echo "$as_me:${as_lineno-$LINENO}: result: no" >&5
$as_echo "no" >&6; }
PKG_CONFIG=""
fi
fi
CONFIG_OBS=
CONFIG_DEPS=
@ -17985,6 +18119,157 @@ if test x${all_targets} = xtrue; then
fi
fi
# AMD debugger API support.
# Check whether --with-amd-dbgapi was given.
if test "${with_amd_dbgapi+set}" = set; then :
withval=$with_amd_dbgapi;
case $withval in
yes | no | auto)
;;
*)
as_fn_error $? "bad value $withval for --with-amd-dbgapi" "$LINENO" 5
;;
esac
else
with_amd_dbgapi=auto
fi
# If the user passes --without-amd-dbgapi but also explicitly enables a target
# that requires amd-dbgapi, it is an error.
if test "$with_amd_dbgapi" = no -a "$gdb_require_amd_dbgapi" = true; then
as_fn_error $? "an explicitly enabled target requires amd-dbgapi, but amd-dbgapi is explicitly disabled" "$LINENO" 5
fi
# Look for amd-dbgapi if:
#
# - a target architecture requiring it has explicitly been enabled, or
# - --enable-targets=all was provided and the user did not explicitly disable
# amd-dbgapi support
if test "$gdb_require_amd_dbgapi" = true \
-o \( "$all_targets" = true -a "$with_amd_dbgapi" != no \); then
# amd-dbgapi version 0.68 is part of ROCm 5.4. There is no guarantee of API
# stability until amd-dbgapi hits 1.0, but for convenience, still check for
# greater or equal that version. It can be handy when testing with a newer
# version of the library.
pkg_failed=no
{ $as_echo "$as_me:${as_lineno-$LINENO}: checking for amd-dbgapi >= 0.68.0" >&5
$as_echo_n "checking for amd-dbgapi >= 0.68.0... " >&6; }
if test -n "$AMD_DBGAPI_CFLAGS"; then
pkg_cv_AMD_DBGAPI_CFLAGS="$AMD_DBGAPI_CFLAGS"
elif test -n "$PKG_CONFIG"; then
if test -n "$PKG_CONFIG" && \
{ { $as_echo "$as_me:${as_lineno-$LINENO}: \$PKG_CONFIG --exists --print-errors \"amd-dbgapi >= 0.68.0\""; } >&5
($PKG_CONFIG --exists --print-errors "amd-dbgapi >= 0.68.0") 2>&5
ac_status=$?
$as_echo "$as_me:${as_lineno-$LINENO}: \$? = $ac_status" >&5
test $ac_status = 0; }; then
pkg_cv_AMD_DBGAPI_CFLAGS=`$PKG_CONFIG --cflags "amd-dbgapi >= 0.68.0" 2>/dev/null`
test "x$?" != "x0" && pkg_failed=yes
else
pkg_failed=yes
fi
else
pkg_failed=untried
fi
if test -n "$AMD_DBGAPI_LIBS"; then
pkg_cv_AMD_DBGAPI_LIBS="$AMD_DBGAPI_LIBS"
elif test -n "$PKG_CONFIG"; then
if test -n "$PKG_CONFIG" && \
{ { $as_echo "$as_me:${as_lineno-$LINENO}: \$PKG_CONFIG --exists --print-errors \"amd-dbgapi >= 0.68.0\""; } >&5
($PKG_CONFIG --exists --print-errors "amd-dbgapi >= 0.68.0") 2>&5
ac_status=$?
$as_echo "$as_me:${as_lineno-$LINENO}: \$? = $ac_status" >&5
test $ac_status = 0; }; then
pkg_cv_AMD_DBGAPI_LIBS=`$PKG_CONFIG --libs "amd-dbgapi >= 0.68.0" 2>/dev/null`
test "x$?" != "x0" && pkg_failed=yes
else
pkg_failed=yes
fi
else
pkg_failed=untried
fi
if test $pkg_failed = no; then
pkg_save_LDFLAGS="$LDFLAGS"
LDFLAGS="$LDFLAGS $pkg_cv_AMD_DBGAPI_LIBS"
cat confdefs.h - <<_ACEOF >conftest.$ac_ext
/* end confdefs.h. */
int
main ()
{
;
return 0;
}
_ACEOF
if ac_fn_c_try_link "$LINENO"; then :
else
pkg_failed=yes
fi
rm -f core conftest.err conftest.$ac_objext \
conftest$ac_exeext conftest.$ac_ext
LDFLAGS=$pkg_save_LDFLAGS
fi
if test $pkg_failed = yes; then
{ $as_echo "$as_me:${as_lineno-$LINENO}: result: no" >&5
$as_echo "no" >&6; }
if $PKG_CONFIG --atleast-pkgconfig-version 0.20; then
_pkg_short_errors_supported=yes
else
_pkg_short_errors_supported=no
fi
if test $_pkg_short_errors_supported = yes; then
AMD_DBGAPI_PKG_ERRORS=`$PKG_CONFIG --short-errors --print-errors --cflags --libs "amd-dbgapi >= 0.68.0" 2>&1`
else
AMD_DBGAPI_PKG_ERRORS=`$PKG_CONFIG --print-errors --cflags --libs "amd-dbgapi >= 0.68.0" 2>&1`
fi
# Put the nasty error message in config.log where it belongs
echo "$AMD_DBGAPI_PKG_ERRORS" >&5
has_amd_dbgapi=no
elif test $pkg_failed = untried; then
{ $as_echo "$as_me:${as_lineno-$LINENO}: result: no" >&5
$as_echo "no" >&6; }
has_amd_dbgapi=no
else
AMD_DBGAPI_CFLAGS=$pkg_cv_AMD_DBGAPI_CFLAGS
AMD_DBGAPI_LIBS=$pkg_cv_AMD_DBGAPI_LIBS
{ $as_echo "$as_me:${as_lineno-$LINENO}: result: yes" >&5
$as_echo "yes" >&6; }
has_amd_dbgapi=yes
fi
if test "$has_amd_dbgapi" = "yes"; then
TARGET_OBS="$TARGET_OBS amd-dbgapi-target.o"
# If --enable-targets=all was provided, use the list of all files depending
# on amd-dbgapi that is hardcoded in the Makefile. Else, the appropriate
# architecture entry in configure.tgt will have added the files to
# gdb_target_obs.
if test "$all_targets" = true; then
TARGET_OBS="$TARGET_OBS \$(ALL_AMD_DBGAPI_TARGET_OBS)"
fi
elif test "$gdb_require_amd_dbgapi" = true -o "$with_amd_dbgapi" = yes; then
# amd-dbgapi was not found and...
#
# - a target requiring it was explicitly enabled, or
# - the user explicitly wants to enable amd-dbgapi
as_fn_error $? "amd-dbgapi is required, but cannot find an appropriate version: $AMD_DBGAPI_PKG_ERRORS" "$LINENO" 5
fi
fi
@ -18087,126 +18372,6 @@ esac
# Handle optional debuginfod support
if test "x$ac_cv_env_PKG_CONFIG_set" != "xset"; then
if test -n "$ac_tool_prefix"; then
# Extract the first word of "${ac_tool_prefix}pkg-config", so it can be a program name with args.
set dummy ${ac_tool_prefix}pkg-config; ac_word=$2
{ $as_echo "$as_me:${as_lineno-$LINENO}: checking for $ac_word" >&5
$as_echo_n "checking for $ac_word... " >&6; }
if ${ac_cv_path_PKG_CONFIG+:} false; then :
$as_echo_n "(cached) " >&6
else
case $PKG_CONFIG in
[\\/]* | ?:[\\/]*)
ac_cv_path_PKG_CONFIG="$PKG_CONFIG" # Let the user override the test with a path.
;;
*)
as_save_IFS=$IFS; IFS=$PATH_SEPARATOR
for as_dir in $PATH
do
IFS=$as_save_IFS
test -z "$as_dir" && as_dir=.
for ac_exec_ext in '' $ac_executable_extensions; do
if as_fn_executable_p "$as_dir/$ac_word$ac_exec_ext"; then
ac_cv_path_PKG_CONFIG="$as_dir/$ac_word$ac_exec_ext"
$as_echo "$as_me:${as_lineno-$LINENO}: found $as_dir/$ac_word$ac_exec_ext" >&5
break 2
fi
done
done
IFS=$as_save_IFS
;;
esac
fi
PKG_CONFIG=$ac_cv_path_PKG_CONFIG
if test -n "$PKG_CONFIG"; then
{ $as_echo "$as_me:${as_lineno-$LINENO}: result: $PKG_CONFIG" >&5
$as_echo "$PKG_CONFIG" >&6; }
else
{ $as_echo "$as_me:${as_lineno-$LINENO}: result: no" >&5
$as_echo "no" >&6; }
fi
fi
if test -z "$ac_cv_path_PKG_CONFIG"; then
ac_pt_PKG_CONFIG=$PKG_CONFIG
# Extract the first word of "pkg-config", so it can be a program name with args.
set dummy pkg-config; ac_word=$2
{ $as_echo "$as_me:${as_lineno-$LINENO}: checking for $ac_word" >&5
$as_echo_n "checking for $ac_word... " >&6; }
if ${ac_cv_path_ac_pt_PKG_CONFIG+:} false; then :
$as_echo_n "(cached) " >&6
else
case $ac_pt_PKG_CONFIG in
[\\/]* | ?:[\\/]*)
ac_cv_path_ac_pt_PKG_CONFIG="$ac_pt_PKG_CONFIG" # Let the user override the test with a path.
;;
*)
as_save_IFS=$IFS; IFS=$PATH_SEPARATOR
for as_dir in $PATH
do
IFS=$as_save_IFS
test -z "$as_dir" && as_dir=.
for ac_exec_ext in '' $ac_executable_extensions; do
if as_fn_executable_p "$as_dir/$ac_word$ac_exec_ext"; then
ac_cv_path_ac_pt_PKG_CONFIG="$as_dir/$ac_word$ac_exec_ext"
$as_echo "$as_me:${as_lineno-$LINENO}: found $as_dir/$ac_word$ac_exec_ext" >&5
break 2
fi
done
done
IFS=$as_save_IFS
;;
esac
fi
ac_pt_PKG_CONFIG=$ac_cv_path_ac_pt_PKG_CONFIG
if test -n "$ac_pt_PKG_CONFIG"; then
{ $as_echo "$as_me:${as_lineno-$LINENO}: result: $ac_pt_PKG_CONFIG" >&5
$as_echo "$ac_pt_PKG_CONFIG" >&6; }
else
{ $as_echo "$as_me:${as_lineno-$LINENO}: result: no" >&5
$as_echo "no" >&6; }
fi
if test "x$ac_pt_PKG_CONFIG" = x; then
PKG_CONFIG=""
else
case $cross_compiling:$ac_tool_warned in
yes:)
{ $as_echo "$as_me:${as_lineno-$LINENO}: WARNING: using cross tools not prefixed with host triplet" >&5
$as_echo "$as_me: WARNING: using cross tools not prefixed with host triplet" >&2;}
ac_tool_warned=yes ;;
esac
PKG_CONFIG=$ac_pt_PKG_CONFIG
fi
else
PKG_CONFIG="$ac_cv_path_PKG_CONFIG"
fi
fi
if test -n "$PKG_CONFIG"; then
_pkg_min_version=0.9.0
{ $as_echo "$as_me:${as_lineno-$LINENO}: checking pkg-config is at least version $_pkg_min_version" >&5
$as_echo_n "checking pkg-config is at least version $_pkg_min_version... " >&6; }
if $PKG_CONFIG --atleast-pkgconfig-version $_pkg_min_version; then
{ $as_echo "$as_me:${as_lineno-$LINENO}: result: yes" >&5
$as_echo "yes" >&6; }
else
{ $as_echo "$as_me:${as_lineno-$LINENO}: result: no" >&5
$as_echo "no" >&6; }
PKG_CONFIG=""
fi
fi
# Handle optional debuginfod support
# Check whether --with-debuginfod was given.

52
gdb/configure.ac
View file

@ -61,6 +61,11 @@ AX_CXX_COMPILE_STDCXX(11, , mandatory)
ZW_CREATE_DEPDIR
ZW_PROG_COMPILER_DEPENDENCIES([CC])
# Since the first call to PKG_CHECK_MODULES may not happen (is guarded by
# a condition), we must call PKG_PROG_PKG_CONFIG explicitly to probe for
# pkg-config.
PKG_PROG_PKG_CONFIG
dnl List of object files and targets accumulated by configure.
CONFIG_OBS=
@ -241,6 +246,53 @@ if test x${all_targets} = xtrue; then
fi
fi
# AMD debugger API support.
AC_ARG_WITH([amd-dbgapi],
[AS_HELP_STRING([--with-amd-dbgapi],
[support for the amd-dbgapi target (yes / no / auto)])],
[GDB_CHECK_YES_NO_AUTO_VAL([$withval], [--with-amd-dbgapi])],
[with_amd_dbgapi=auto])
# If the user passes --without-amd-dbgapi but also explicitly enables a target
# that requires amd-dbgapi, it is an error.
if test "$with_amd_dbgapi" = no -a "$gdb_require_amd_dbgapi" = true; then
AC_MSG_ERROR([an explicitly enabled target requires amd-dbgapi, but amd-dbgapi is explicitly disabled])
fi
# Look for amd-dbgapi if:
#
# - a target architecture requiring it has explicitly been enabled, or
# - --enable-targets=all was provided and the user did not explicitly disable
# amd-dbgapi support
if test "$gdb_require_amd_dbgapi" = true \
-o \( "$all_targets" = true -a "$with_amd_dbgapi" != no \); then
# amd-dbgapi version 0.68 is part of ROCm 5.4. There is no guarantee of API
# stability until amd-dbgapi hits 1.0, but for convenience, still check for
# greater or equal that version. It can be handy when testing with a newer
# version of the library.
PKG_CHECK_MODULES([AMD_DBGAPI], [amd-dbgapi >= 0.68.0],
[has_amd_dbgapi=yes], [has_amd_dbgapi=no])
if test "$has_amd_dbgapi" = "yes"; then
TARGET_OBS="$TARGET_OBS amd-dbgapi-target.o"
# If --enable-targets=all was provided, use the list of all files depending
# on amd-dbgapi that is hardcoded in the Makefile. Else, the appropriate
# architecture entry in configure.tgt will have added the files to
# gdb_target_obs.
if test "$all_targets" = true; then
TARGET_OBS="$TARGET_OBS \$(ALL_AMD_DBGAPI_TARGET_OBS)"
fi
elif test "$gdb_require_amd_dbgapi" = true -o "$with_amd_dbgapi" = yes; then
# amd-dbgapi was not found and...
#
# - a target requiring it was explicitly enabled, or
# - the user explicitly wants to enable amd-dbgapi
AC_MSG_ERROR([amd-dbgapi is required, but cannot find an appropriate version: $AMD_DBGAPI_PKG_ERRORS])
fi
fi
AC_SUBST(TARGET_OBS)
AC_SUBST(HAVE_NATIVE_GCORE_TARGET)

23
gdb/configure.tgt
View file

@ -2,13 +2,20 @@
# invoked from the autoconf generated configure script.
# This file sets the following shell variables:
# gdb_target_obs target-specific object files to use
# gdb_sim simulator library for target
# gdb_osabi default OS ABI to use with target
# gdb_have_gcore set to "true"/"false" if this target can run gcore
# gdb_target_obs target-specific object files to use
# gdb_sim simulator library for target
# gdb_osabi default OS ABI to use with target
# gdb_have_gcore set to "true"/"false" if this target can run gcore
# gdb_require_amd_dbgapi set to "true" if this target requires the amd-dbgapi
# target
# NOTE: Every file added to a gdb_target_obs variable for any target here
# must also be added to either ALL_TARGET_OBS or ALL_64_TARGET_OBS
# must also be added to either:
#
# - ALL_TARGET_OBS
# - ALL_64_TARGET_OBS
# - ALL_AMD_DBGAPI_TARGET_OBS
#
# in Makefile.in!
case $targ in
@ -161,6 +168,12 @@ alpha*-*-openbsd*)
alpha-netbsd-tdep.o alpha-obsd-tdep.o netbsd-tdep.o"
;;
amdgcn*-*-*)
# Target: AMDGPU
gdb_require_amd_dbgapi=true
gdb_target_obs="amdgpu-tdep.o solib-rocm.o"
;;
am33_2.0*-*-linux*)
# Target: Matsushita mn10300 (AM33) running Linux
gdb_target_obs="mn10300-tdep.o mn10300-linux-tdep.o linux-tdep.o \

291
gdb/doc/gdb.texinfo
View file

@ -7026,6 +7026,8 @@ signal happened. @value{GDBN} alerts you to the context switch with a
message such as @samp{[Switching to Thread @var{n}]} to identify the
thread.
@anchor{set scheduler-locking}
On some OSes, you can modify @value{GDBN}'s default behavior by
locking the OS scheduler to allow only a single thread to run.
@ -25882,6 +25884,7 @@ all uses of @value{GDBN} with the architecture, both native and cross.
* Nios II::
* Sparc64::
* S12Z::
* AMD GPU:: @acronym{AMD GPU} architectures
@end menu
@node AArch64
@ -26370,6 +26373,254 @@ This command displays the current value of the microprocessor's
BDCCSR register.
@end table
@node AMD GPU
@subsection @acronym{AMD GPU}
@cindex @acronym{AMD GPU} support
@value{GDBN} supports debugging programs offloaded to @acronym{AMD GPU} devices
using the @url{https://docs.amd.com/, @acronym{AMD ROCm}} platform.
@value{GDBN} presents host threads alongside GPU wavefronts, allowing debugging
both the host and device parts of the program simultaneously.
@subsubsection @acronym{AMD GPU} Architectures
The list of @acronym{AMD GPU} architectures supported by @value{GDBN} depends
on the version of the AMD Debugger API library used. See its
@uref{https://docs.amd.com/bundle/ROCDebugger_User_and_API, documentation} for
more details.
@subsubsection @acronym{AMD GPU} Device Driver and @acronym{AMD ROCm} Runtime
@value{GDBN} requires a compatible @acronym{AMD GPU} device driver to
be installed. A warning message is displayed if either the device
driver version or the version of the debug support it implements is
unsupported. @value{GDBN} will continue to function except no
@acronym{AMD GPU} debugging will be possible.
@value{GDBN} requires each agent to have compatible firmware installed
by the device driver. A warning message is displayed if unsupported
firmware is detected. @value{GDBN} will continue to function except
no @acronym{AMD GPU} debugging will be possible on the agent.
@value{GDBN} requires a compatible @acronym{AMD ROCm} runtime to be
loaded in order to detect @acronym{AMD GPU} code objects and
wavefronts. A warning message is displayed if an unsupported
@acronym{AMD ROCm} runtime is detected, or there is an error or
restriction that prevents debugging. @value{GDBN} will continue to
function except no @acronym{AMD GPU} debugging will be possible.
@subsubsection @acronym{AMD GPU} Wavefronts
@cindex wavefronts
An @acronym{AMD GPU} wavefront is represented in @value{GDBN} as a
thread.
Note that some @acronym{AMD GPU} architectures may have restrictions
on providing information about @acronym{AMD GPU} wavefronts created
when @value{GDBN} is not attached (@pxref{AMD GPU Attaching
Restrictions, , @acronym{AMD GPU} Attaching Restrictions}).
When scheduler-locking is in effect (@pxref{set scheduler-locking}),
new wavefronts created by the resumed thread (either CPU thread or GPU
wavefront) are held in the halt state.
@subsubsection @acronym{AMD GPU} Code Objects
The @samp{info sharedlibrary} command will show the @acronym{AMD GPU}
code objects as file or memory URIs, together with the host's shared
libraries. For example:
@smallexample
(@value{GDBP}) info sharedlibrary
From To Syms Read Shared Object Library
0x1111 0x2222 Yes (*) /lib64/ld-linux-x86-64.so.2
...
0x3333 0x4444 Yes (*) /opt/rocm-4.5.0/.../libamd_comgr.so
0x5555 0x6666 Yes (*) /lib/x86_64-linux-gnu/libtinfo.so.5
0x7777 0x8888 Yes file:///tmp/a.out#offset=6477&size=10832
0x9999 0xaaaa Yes (*) memory://95557/mem#offset=0x1234&size=100
(*): Shared library is missing debugging information.
(@value{GDBP})
@end smallexample
For a @samp{file} URI, the path portion is the file on disk containing
the code object. The @var{offset} parameter is a 0-based offset in
this file, to the start of the code object. If omitted, it defaults to
0. The @var{size} parameter is the size of the code object in bytes.
If omitted, it defaults to the size of the file.
For a @samp{memory} URI, the path portion is the process id of the
process owning the memory containing the code object. The @var{offset}
parameter is the memory address where the code object is found, and
the @var{size} parameter is its size in bytes.
@acronym{AMD GPU} code objects are loaded into each @acronym{AMD GPU}
device separately. The @samp{info sharedlibrary} command may
therefore show the same code object loaded multiple times. As a
consequence, setting a breakpoint in @acronym{AMD GPU} code will
result in multiple breakpoint locations if there are multiple
@acronym{AMD GPU} devices.
@subsubsection @acronym{AMD GPU} Entity Target Identifiers and Convenience Variables
The @acronym{AMD GPU} entities have the following target identifier formats:
@table @asis
@item Thread Target ID
The @acronym{AMD GPU} thread target identifier (@var{systag}) string has the
following format:
@smallexample
AMDGPU Wave @var{agent-id}:@var{queue-id}:@var{dispatch-id}:@var{wave-id} (@var{work-group-x},@var{work-group-y},@var{work-group-z})/@var{work-group-thread-index}
@end smallexample
@end table
@anchor{AMD GPU Signals}
@subsubsection @acronym{AMD GPU} Signals
For @acronym{AMD GPU} wavefronts, @value{GDBN} maps target conditions to stop
signals in the following way:
@table @code
@item SIGILL
Execution of an illegal instruction.
@item SIGTRAP
Execution of a @code{S_TRAP} instruction other than:
@itemize @bullet{}
@item
@code{S_TRAP 1} which is used by @value{GDBN} to insert breakpoints.
@item
@code{S_TRAP 2} which raises @code{SIGABRT}.
@end itemize
@item SIGABRT
Execution of a @code{S_TRAP 2} instruction.
@item SIGFPE
Execution of a floating point or integer instruction detects a
condition that is enabled to raise a signal. The conditions include:
@itemize @bullet{}
@item
Floating point operation is invalid.
@item
Floating point operation had subnormal input that was rounded to zero.
@item
Floating point operation performed a division by zero.
@item
Floating point operation produced an overflow result. The result was
rounded to infinity.
@item
Floating point operation produced an underflow result. A subnormal
result was rounded to zero.
@item
Floating point operation produced an inexact result.
@item
Integer operation performed a division by zero.
@end itemize
By default, these conditions are not enabled to raise signals. The
@samp{set $mode} command can be used to change the @acronym{AMD GPU}
wavefront's register that has bits controlling which conditions are
enabled to raise signals. The @samp{print $trapsts} command can be
used to inspect which conditions have been detected even if they are
not enabled to raise a signal.
@item SIGBUS
Execution of an instruction that accessed global memory using an
address that is outside the virtual address range.
@item SIGSEGV
Execution of an instruction that accessed a global memory page that is
either not mapped or accessed with incompatible permissions.
@end table
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} Logging
The @samp{set debug amd-dbgapi} command can be used
to enable diagnostic messages in the @samp{amd-dbgapi} target. The
@samp{show debug amd-dbgapi} command displays the current setting.
@xref{set debug amd-dbgapi}.
The @samp{set debug amd-dbgapi-lib log-level @var{level}} command can be used
to enable diagnostic messages from the @samp{amd-dbgapi} library (which
@value{GDBN} uses under the hood). The @samp{show debug amd-dbgapi-lib
log-level} command displays the current @samp{amd-dbgapi} library log level.
@xref{set debug amd-dbgapi-lib}.
@subsubsection @acronym{AMD GPU} Restrictions
@enumerate
@item
When in non-stop mode, wavefronts may not hit breakpoints inserted
while not stopped, nor see memory updates made while not stopped,
until the wavefront is next stopped. Memory updated by non-stopped
wavefronts may not be visible until the wavefront is next stopped.
@item The HIP runtime performs deferred code object loading by default.
@acronym{AMD GPU} code objects are not loaded until the first kernel is
launched. Before then, all breakpoints have to be set as pending breakpoints.
If source line positions are used that only correspond to source lines in
unloaded code objects, then @value{GDBN} may not set pending breakpoints, and
instead set breakpoints on the next following source line that maps to host
code. This can result in unexpected breakpoint hits being reported. When the
code object containing the source lines is loaded, the incorrect breakpoints
will be removed and replaced by the correct ones. This problem can be avoided
by only setting breakpoints in unloaded code objects using symbol or function
names.
Setting the @code{HIP_ENABLE_DEFERRED_LOADING} environment variable to @code{0}
can be used to disable deferred code object loading by the HIP runtime. This
ensures all code objects will be loaded when the inferior reaches the beginning
of the @code{main} function.
@item
If no CPU thread is running, then @samp{Ctrl-C} is not able to stop
@acronym{AMD GPU} threads. This can happen for example if you enable
@code{scheduler-locking} after the whole program stopped, and then resume an
@acronym{AMD GPU} thread. The only way to unblock the situation is to kill the
@value{GDBN} process.
@anchor{AMD GPU Attaching Restrictions}
@item
By default, for some architectures, the @acronym{AMD GPU} device driver causes
all @acronym{AMD GPU} wavefronts created when @value{GDBN} is not attached to
be unable to report the dispatch associated with the wavefront, or the
wavefront's work-group position. The @samp{info threads} command will display
this missing information with a @samp{?}.
This does not affect wavefronts created while @value{GDBN} is attached which
are always capable of reporting this information.
If the @env{HSA_ENABLE_DEBUG} environment variable is set to @samp{1} when the
@acronym{AMD ROCm} runtime is initialized, then this information will be
available for all architectures even for wavefronts created when @value{GDBN}
was not attached.
@end enumerate
@node Controlling GDB
@chapter Controlling @value{GDBN}
@ -27623,6 +27874,46 @@ module.
@item show debug aix-thread
Show the current state of AIX thread debugging info display.
@cindex AMD GPU debugging info
@anchor{set debug amd-dbgapi-lib}
@item set debug amd-dbgapi-lib
@itemx show debug amd-dbgapi-lib
The @code{set debug amd-dbgapi-lib log-level @var{level}} command can be used
to enable diagnostic messages from the @samp{amd-dbgapi} library, where
@var{level} can be:
@table @code
@item off
no logging is enabled
@item error
fatal errors are reported
@item warning
fatal errors and warnings are reported
@item info
fatal errors, warnings, and info messages are reported
@item verbose
all messages are reported
@end table
The @code{show debug amd-dbgapi-lib log-level} command displays the current
@acronym{amd-dbgapi} library log level.
@anchor{set debug amd-dbgapi}
@item set debug amd-dbgapi
@itemx show debug amd-dbgapi
The @samp{set debug amd-dbgapi} command can be used
to enable diagnostic messages in the @samp{amd-dbgapi} target. The
@samp{show debug amd-dbgapi} command displays the current setting.
@xref{set debug amd-dbgapi}.
@item set debug check-physname
@cindex physname
Check the results of the ``physname'' computation. When reading DWARF

3
gdb/regcache.c
View file

@ -1915,7 +1915,8 @@ cooked_read_test (struct gdbarch *gdbarch)
{
auto bfd_arch = gdbarch_bfd_arch_info (gdbarch)->arch;
if (bfd_arch == bfd_arch_frv || bfd_arch == bfd_arch_h8300
if (bfd_arch == bfd_arch_amdgcn
|| bfd_arch == bfd_arch_frv || bfd_arch == bfd_arch_h8300
|| bfd_arch == bfd_arch_m32c || bfd_arch == bfd_arch_sh
|| bfd_arch == bfd_arch_alpha || bfd_arch == bfd_arch_v850
|| bfd_arch == bfd_arch_msp430 || bfd_arch == bfd_arch_mep

679
gdb/solib-rocm.c Normal file
View file

@ -0,0 +1,679 @@
/* Handle ROCm Code Objects for GDB, the GNU Debugger.
Copyright (C) 2019-2022 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 "defs.h"
#include "amd-dbgapi-target.h"
#include "amdgpu-tdep.h"
#include "arch-utils.h"
#include "elf-bfd.h"
#include "elf/amdgpu.h"
#include "gdbsupport/fileio.h"
#include "inferior.h"
#include "observable.h"
#include "solib.h"
#include "solib-svr4.h"
#include "solist.h"
#include "symfile.h"
/* ROCm-specific inferior data. */
struct solib_info
{
/* List of code objects loaded into the inferior. */
so_list *solib_list;
};
/* Per-inferior data key. */
static const registry<inferior>::key<solib_info> rocm_solib_data;
static target_so_ops rocm_solib_ops;
/* Free the solib linked list. */
static void
rocm_free_solib_list (struct solib_info *info)
{
while (info->solib_list != nullptr)
{
struct so_list *next = info->solib_list->next;
free_so (info->solib_list);
info->solib_list = next;
}
info->solib_list = nullptr;
}
/* Fetch the solib_info data for INF. */
static struct solib_info *
get_solib_info (inferior *inf)
{
solib_info *info = rocm_solib_data.get (inf);
if (info == nullptr)
info = rocm_solib_data.emplace (inf);
return info;
}
/* Relocate section addresses. */
static void
rocm_solib_relocate_section_addresses (struct so_list *so,
struct target_section *sec)
{
if (!is_amdgpu_arch (gdbarch_from_bfd (so->abfd)))
{
svr4_so_ops.relocate_section_addresses (so, sec);
return;
}
lm_info_svr4 *li = (lm_info_svr4 *) so->lm_info;
sec->addr = sec->addr + li->l_addr;
sec->endaddr = sec->endaddr + li->l_addr;
}
static void rocm_update_solib_list ();
static void
rocm_solib_handle_event ()
{
/* Since we sit on top of svr4_so_ops, we might get called following an event
concerning host libraries. We must therefore forward the call. If the
event was for a ROCm code object, it will be a no-op. On the other hand,
if the event was for host libraries, rocm_update_solib_list will be
essentially be a no-op (it will reload the same code object list as was
previously loaded). */
svr4_so_ops.handle_event ();
rocm_update_solib_list ();
}
/* Make a deep copy of the solib linked list. */
static so_list *
rocm_solib_copy_list (const so_list *src)
{
struct so_list *dst = nullptr;
struct so_list **link = &dst;
while (src != nullptr)
{
struct so_list *newobj;
newobj = XNEW (struct so_list);
memcpy (newobj, src, sizeof (struct so_list));
lm_info_svr4 *src_li = (lm_info_svr4 *) src->lm_info;
newobj->lm_info = new lm_info_svr4 (*src_li);
newobj->next = nullptr;
*link = newobj;
link = &newobj->next;
src = src->next;
}
return dst;
}
/* Build a list of `struct so_list' objects describing the shared
objects currently loaded in the inferior. */
static struct so_list *
rocm_solib_current_sos ()
{
/* First, retrieve the host-side shared library list. */
so_list *head = svr4_so_ops.current_sos ();
/* Then, the device-side shared library list. */
so_list *list = get_solib_info (current_inferior ())->solib_list;
if (list == nullptr)
return head;
list = rocm_solib_copy_list (list);
if (head == nullptr)
return list;
/* Append our libraries to the end of the list. */
so_list *tail;
for (tail = head; tail->next; tail = tail->next)
/* Nothing. */;
tail->next = list;
return head;
}
namespace {
/* Interface to interact with a ROCm code object stream. */
struct rocm_code_object_stream
{
DISABLE_COPY_AND_ASSIGN (rocm_code_object_stream);
/* Copy SIZE bytes from the underlying objfile storage starting at OFFSET
into the user provided buffer BUF.
Return the number of bytes actually copied (might be inferior to SIZE if
the end of the stream is reached). */
virtual file_ptr read (void *buf, file_ptr size, file_ptr offset) = 0;
/* Retrieve file information in SB.
Return 0 on success. On failure, set the appropriate bfd error number
(using bfd_set_error) and return -1. */
int stat (struct stat *sb);
virtual ~rocm_code_object_stream () = default;
protected:
rocm_code_object_stream () = default;
/* Return the size of the object file, or -1 if the size cannot be
determined.
This is a helper function for stat. */
virtual LONGEST size () = 0;
};
int
rocm_code_object_stream::stat (struct stat *sb)
{
const LONGEST size = this->size ();
if (size == -1)
return -1;
memset (sb, '\0', sizeof (struct stat));
sb->st_size = size;
return 0;
}
/* Interface to a ROCm object stream which is embedded in an ELF file
accessible to the debugger. */
struct rocm_code_object_stream_file final : rocm_code_object_stream
{
DISABLE_COPY_AND_ASSIGN (rocm_code_object_stream_file);
rocm_code_object_stream_file (int fd, ULONGEST offset, ULONGEST size);
file_ptr read (void *buf, file_ptr size, file_ptr offset) override;
LONGEST size () override;
~rocm_code_object_stream_file () override;
protected:
/* The target file descriptor for this stream. */
int m_fd;
/* The offset of the ELF file image in the target file. */
ULONGEST m_offset;
/* The size of the ELF file image. The value 0 means that it was
unspecified in the URI descriptor. */
ULONGEST m_size;
};
rocm_code_object_stream_file::rocm_code_object_stream_file
(int fd, ULONGEST offset, ULONGEST size)
: m_fd (fd), m_offset (offset), m_size (size)
{
}
file_ptr
rocm_code_object_stream_file::read (void *buf, file_ptr size,
file_ptr offset)
{
fileio_error target_errno;
file_ptr nbytes = 0;
while (size > 0)
{
QUIT;
file_ptr bytes_read
= target_fileio_pread (m_fd, static_cast<gdb_byte *> (buf) + nbytes,
size, m_offset + offset + nbytes,
&target_errno);
if (bytes_read == 0)
break;
if (bytes_read < 0)
{
errno = fileio_error_to_host (target_errno);
bfd_set_error (bfd_error_system_call);
return -1;
}
nbytes += bytes_read;
size -= bytes_read;
}
return nbytes;
}
LONGEST
rocm_code_object_stream_file::size ()
{
if (m_size == 0)
{
fileio_error target_errno;
struct stat stat;
if (target_fileio_fstat (m_fd, &stat, &target_errno) < 0)
{
errno = fileio_error_to_host (target_errno);
bfd_set_error (bfd_error_system_call);
return -1;
}
/* Check that the offset is valid. */
if (m_offset >= stat.st_size)
{
bfd_set_error (bfd_error_bad_value);
return -1;
}
m_size = stat.st_size - m_offset;
}
return m_size;
}
rocm_code_object_stream_file::~rocm_code_object_stream_file ()
{
fileio_error target_errno;
target_fileio_close (m_fd, &target_errno);
}
/* Interface to a code object which lives in the inferior's memory. */
struct rocm_code_object_stream_memory final : public rocm_code_object_stream
{
DISABLE_COPY_AND_ASSIGN (rocm_code_object_stream_memory);
rocm_code_object_stream_memory (gdb::byte_vector buffer);
file_ptr read (void *buf, file_ptr size, file_ptr offset) override;
protected:
/* Snapshot of the original ELF image taken during load. This is done to
support the situation where an inferior uses an in-memory image, and
releases or re-uses this memory before GDB is done using it. */
gdb::byte_vector m_objfile_image;
LONGEST size () override
{
return m_objfile_image.size ();
}
};
rocm_code_object_stream_memory::rocm_code_object_stream_memory
(gdb::byte_vector buffer)
: m_objfile_image (std::move (buffer))
{
}
file_ptr
rocm_code_object_stream_memory::read (void *buf, file_ptr size,
file_ptr offset)
{
if (size > m_objfile_image.size () - offset)
size = m_objfile_image.size () - offset;
memcpy (buf, m_objfile_image.data () + offset, size);
return size;
}
} /* anonymous namespace */
static void *
rocm_bfd_iovec_open (bfd *abfd, void *inferior_void)
{
gdb::string_view uri (bfd_get_filename (abfd));
gdb::string_view protocol_delim = "://";
size_t protocol_end = uri.find (protocol_delim);
std::string protocol = gdb::to_string (uri.substr (0, protocol_end));
protocol_end += protocol_delim.length ();
std::transform (protocol.begin (), protocol.end (), protocol.begin (),
[] (unsigned char c) { return std::tolower (c); });
gdb::string_view path;
size_t path_end = uri.find_first_of ("#?", protocol_end);
if (path_end != std::string::npos)
path = uri.substr (protocol_end, path_end++ - protocol_end);
else
path = uri.substr (protocol_end);
/* %-decode the string. */
std::string decoded_path;
decoded_path.reserve (path.length ());
for (size_t i = 0; i < path.length (); ++i)
if (path[i] == '%'
&& i < path.length () - 2
&& std::isxdigit (path[i + 1])
&& std::isxdigit (path[i + 2]))
{
gdb::string_view hex_digits = path.substr (i + 1, 2);
decoded_path += std::stoi (gdb::to_string (hex_digits), 0, 16);
i += 2;
}
else
decoded_path += path[i];
/* Tokenize the query/fragment. */
std::vector<gdb::string_view> tokens;
size_t pos, last = path_end;
while ((pos = uri.find ('&', last)) != std::string::npos)
{
tokens.emplace_back (uri.substr (last, pos - last));
last = pos + 1;
}
if (last != std::string::npos)
tokens.emplace_back (uri.substr (last));
/* Create a tag-value map from the tokenized query/fragment. */
std::unordered_map<gdb::string_view, gdb::string_view,
gdb::string_view_hash> params;
for (gdb::string_view token : tokens)
{
size_t delim = token.find ('=');
if (delim != std::string::npos)
{
gdb::string_view tag = token.substr (0, delim);
gdb::string_view val = token.substr (delim + 1);
params.emplace (tag, val);
}
}
try
{
ULONGEST offset = 0;
ULONGEST size = 0;
inferior *inferior = static_cast<struct inferior *> (inferior_void);
auto try_strtoulst = [] (gdb::string_view v)
{
errno = 0;
ULONGEST value = strtoulst (v.data (), nullptr, 0);
if (errno != 0)
{
/* The actual message doesn't matter, the exception is caught
below, transformed in a BFD error, and the message is lost. */
error (_("Failed to parse integer."));
}
return value;
};
auto offset_it = params.find ("offset");
if (offset_it != params.end ())
offset = try_strtoulst (offset_it->second);
auto size_it = params.find ("size");
if (size_it != params.end ())
{
size = try_strtoulst (size_it->second);
if (size == 0)
error (_("Invalid size value"));
}
if (protocol == "file")
{
fileio_error target_errno;
int fd
= target_fileio_open (static_cast<struct inferior *> (inferior),
decoded_path.c_str (), FILEIO_O_RDONLY,
false, 0, &target_errno);
if (fd == -1)
{
errno = fileio_error_to_host (target_errno);
bfd_set_error (bfd_error_system_call);
return nullptr;
}
return new rocm_code_object_stream_file (fd, offset, size);
}
if (protocol == "memory")
{
ULONGEST pid = try_strtoulst (path);
if (pid != inferior->pid)
{
warning (_("`%s': code object is from another inferior"),
gdb::to_string (uri).c_str ());
bfd_set_error (bfd_error_bad_value);
return nullptr;
}
gdb::byte_vector buffer (size);
if (target_read_memory (offset, buffer.data (), size) != 0)
{
warning (_("Failed to copy the code object from the inferior"));
bfd_set_error (bfd_error_bad_value);
return nullptr;
}
return new rocm_code_object_stream_memory (std::move (buffer));
}
warning (_("`%s': protocol not supported: %s"),
gdb::to_string (uri).c_str (), protocol.c_str ());
bfd_set_error (bfd_error_bad_value);
return nullptr;
}
catch (const gdb_exception_quit &ex)
{
set_quit_flag ();
bfd_set_error (bfd_error_bad_value);
return nullptr;
}
catch (const gdb_exception &ex)
{
bfd_set_error (bfd_error_bad_value);
return nullptr;
}
}
static int
rocm_bfd_iovec_close (bfd *nbfd, void *data)
{
delete static_cast<rocm_code_object_stream *> (data);
return 0;
}
static file_ptr
rocm_bfd_iovec_pread (bfd *abfd, void *data, void *buf, file_ptr size,
file_ptr offset)
{
return static_cast<rocm_code_object_stream *> (data)->read (buf, size,
offset);
}
static int
rocm_bfd_iovec_stat (bfd *abfd, void *data, struct stat *sb)
{
return static_cast<rocm_code_object_stream *> (data)->stat (sb);
}
static gdb_bfd_ref_ptr
rocm_solib_bfd_open (const char *pathname)
{
/* Handle regular files with SVR4 open. */
if (strstr (pathname, "://") == nullptr)
return svr4_so_ops.bfd_open (pathname);
gdb_bfd_ref_ptr abfd
= gdb_bfd_openr_iovec (pathname, "elf64-amdgcn", rocm_bfd_iovec_open,
current_inferior (), rocm_bfd_iovec_pread,
rocm_bfd_iovec_close, rocm_bfd_iovec_stat);
if (abfd == nullptr)
error (_("Could not open `%s' as an executable file: %s"), pathname,
bfd_errmsg (bfd_get_error ()));
/* Check bfd format. */
if (!bfd_check_format (abfd.get (), bfd_object))
error (_("`%s': not in executable format: %s"),
bfd_get_filename (abfd.get ()), bfd_errmsg (bfd_get_error ()));
unsigned char osabi = elf_elfheader (abfd)->e_ident[EI_OSABI];
unsigned char osabiversion = elf_elfheader (abfd)->e_ident[EI_ABIVERSION];
/* Check that the code object is using the HSA OS ABI. */
if (osabi != ELFOSABI_AMDGPU_HSA)
error (_("`%s': ELF file OS ABI is not supported (%d)."),
bfd_get_filename (abfd.get ()), osabi);
/* We support HSA code objects V3 and greater. */
if (osabiversion < ELFABIVERSION_AMDGPU_HSA_V3)
error (_("`%s': ELF file HSA OS ABI version is not supported (%d)."),
bfd_get_filename (abfd.get ()), osabiversion);
return abfd;
}
static void
rocm_solib_create_inferior_hook (int from_tty)
{
rocm_free_solib_list (get_solib_info (current_inferior ()));
svr4_so_ops.solib_create_inferior_hook (from_tty);
}
static void
rocm_update_solib_list ()
{
inferior *inf = current_inferior ();
amd_dbgapi_process_id_t process_id = get_amd_dbgapi_process_id (inf);
if (process_id.handle == AMD_DBGAPI_PROCESS_NONE.handle)
return;
solib_info *info = get_solib_info (inf);
rocm_free_solib_list (info);
struct so_list **link = &info->solib_list;
amd_dbgapi_code_object_id_t *code_object_list;
size_t count;
amd_dbgapi_status_t status
= amd_dbgapi_process_code_object_list (process_id, &count,
&code_object_list, nullptr);
if (status != AMD_DBGAPI_STATUS_SUCCESS)
{
warning (_("amd_dbgapi_process_code_object_list failed (%s)"),
get_status_string (status));
return;
}
for (size_t i = 0; i < count; ++i)
{
CORE_ADDR l_addr;
char *uri_bytes;
status = amd_dbgapi_code_object_get_info
(code_object_list[i], AMD_DBGAPI_CODE_OBJECT_INFO_LOAD_ADDRESS,
sizeof (l_addr), &l_addr);
if (status != AMD_DBGAPI_STATUS_SUCCESS)
continue;
status = amd_dbgapi_code_object_get_info
(code_object_list[i], AMD_DBGAPI_CODE_OBJECT_INFO_URI_NAME,
sizeof (uri_bytes), &uri_bytes);
if (status != AMD_DBGAPI_STATUS_SUCCESS)
continue;
struct so_list *so = XCNEW (struct so_list);
lm_info_svr4 *li = new lm_info_svr4;
li->l_addr = l_addr;
so->lm_info = li;
strncpy (so->so_name, uri_bytes, sizeof (so->so_name));
so->so_name[sizeof (so->so_name) - 1] = '\0';
xfree (uri_bytes);
/* Make so_original_name unique so that code objects with the same URI
but different load addresses are seen by gdb core as different shared
objects. */
xsnprintf (so->so_original_name, sizeof (so->so_original_name),
"code_object_%ld", code_object_list[i].handle);
so->next = nullptr;
*link = so;
link = &so->next;
}
xfree (code_object_list);
if (rocm_solib_ops.current_sos == NULL)
{
/* Override what we need to. */
rocm_solib_ops = svr4_so_ops;
rocm_solib_ops.current_sos = rocm_solib_current_sos;
rocm_solib_ops.solib_create_inferior_hook
= rocm_solib_create_inferior_hook;
rocm_solib_ops.bfd_open = rocm_solib_bfd_open;
rocm_solib_ops.relocate_section_addresses
= rocm_solib_relocate_section_addresses;
rocm_solib_ops.handle_event = rocm_solib_handle_event;
/* Engage the ROCm so_ops. */
set_gdbarch_so_ops (current_inferior ()->gdbarch, &rocm_solib_ops);
}
}
static void
rocm_solib_target_inferior_created (inferior *inf)
{
rocm_free_solib_list (get_solib_info (inf));
rocm_update_solib_list ();
/* Force GDB to reload the solibs. */
current_inferior ()->pspace->clear_solib_cache ();
solib_add (nullptr, 0, auto_solib_add);
}
/* -Wmissing-prototypes */
extern initialize_file_ftype _initialize_rocm_solib;
void
_initialize_rocm_solib ()
{
/* The dependency on the amd-dbgapi exists because solib-rocm's
inferior_created observer needs amd-dbgapi to have attached the process,
which happens in amd_dbgapi_target's inferior_created observer. */
gdb::observers::inferior_created.attach
(rocm_solib_target_inferior_created,
"solib-rocm",
{ &get_amd_dbgapi_target_inferior_created_observer_token () });
}

48
gdb/testsuite/gdb.rocm/simple.cpp Normal file
View file

@ -0,0 +1,48 @@
/* This testcase is part of GDB, the GNU debugger.
Copyright 2022 Free Software Foundation, Inc.
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"
#include <cassert>
__global__ void
do_an_addition (int a, int b, int *out)
{
*out = a + b;
}
int
main ()
{
int *result_ptr, result;
/* Allocate memory for the device to write the result to. */
hipError_t error = hipMalloc (&result_ptr, sizeof (int));
assert (error == hipSuccess);
/* Run `do_an_addition` on one workgroup containing one work item. */
do_an_addition<<<dim3(1), dim3(1), 0, 0>>> (1, 2, result_ptr);
/* Copy result from device to host. Note that this acts as a synchronization
point, waiting for the kernel dispatch to complete. */
error = hipMemcpyDtoH (&result, result_ptr, sizeof (int));
assert (error == hipSuccess);
printf ("result is %d\n", result);
assert (result == 3);
return 0;
}

52
gdb/testsuite/gdb.rocm/simple.exp Normal file
View file

@ -0,0 +1,52 @@
# Copyright 2022 Free Software Foundation, Inc.
# 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/>.
# A simple AMD GPU debugging smoke test. Run to a breakpoint in device code,
# then continue until the end of the program.
load_lib rocm.exp
standard_testfile .cpp
if [skip_hipcc_tests] {
verbose "skipping hip test: ${testfile}"
return
}
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 "with breakpoint pending on -- break do_an_addition" \
"Breakpoint $::decimal \\(do_an_addition\\) pending."
gdb_test "continue" \
"Thread $::decimal hit Breakpoint $::decimal, do_an_addition .*"
gdb_test "continue" \
"Inferior 1 .* exited normally.*" \
"continue to end"
}
}
do_test

38
gdb/testsuite/lib/future.exp
View file

@ -121,6 +121,19 @@ proc gdb_find_rustc {} {
return $rustc
}
proc gdb_find_hipcc {} {
global tool_root_dir
if {![is_remote host]} {
set hipcc [lookfor_file $tool_root_dir hipcc]
if {$hipcc == ""} {
set hipcc [lookfor_file /opt/rocm/bin hipcc]
}
} else {
set hipcc ""
}
return $hipcc
}
proc gdb_find_ldd {} {
global LDD_FOR_TARGET
if [info exists LDD_FOR_TARGET] {
@ -290,6 +303,18 @@ proc gdb_default_target_compile_1 {source destfile type options} {
}
}
if { $i == "hip" } {
set compiler_type "hip"
if {[board_info $dest exists hipflags]} {
append add_flags " [target_info hipflags]"
}
if {[board_info $dest exists hipcompiler]} {
set compiler [target_info hipcompiler]
} else {
set compiler [find_hipcc]
}
}
if {[regexp "^dest=" $i]} {
regsub "^dest=" $i "" tmp
if {[board_info $tmp exists name]} {
@ -352,6 +377,7 @@ proc gdb_default_target_compile_1 {source destfile type options} {
global GO_FOR_TARGET
global GO_LD_FOR_TARGET
global RUSTC_FOR_TARGET
global HIPCC_FOR_TARGET
if {[info exists GNATMAKE_FOR_TARGET]} {
if { $compiler_type == "ada" } {
@ -398,6 +424,12 @@ proc gdb_default_target_compile_1 {source destfile type options} {
}
}
if {[info exists HIPCC_FOR_TARGET]} {
if {$compiler_type == "hip"} {
set compiler $HIPCC_FOR_TARGET
}
}
if { $type == "executable" && $linker != "" } {
set compiler $linker
}
@ -687,6 +719,12 @@ if {[info procs find_rustc] == ""} {
gdb_note [join [list $note_prefix "Rust" $note_suffix] ""]
}
if {[info procs find_hipcc] == ""} {
rename gdb_find_hipcc find_hipcc
set use_gdb_compile(hip) 1
gdb_note [join [list $note_prefix "HIP" $note_suffix] ""]
}
# If dejagnu's default_target_compile is missing support for any language,
# override it.
if { [array size use_gdb_compile] != 0 } {

7
gdb/testsuite/lib/gdb.exp
View file

@ -4867,6 +4867,13 @@ proc gdb_compile {source dest type options} {
lappend new_options "early_flags=-fno-stack-protector"
}
# hipcc defaults to -O2, so add -O0 to early flags for the hip language.
# If "optimize" is also requested, another -O flag (e.g. -O2) will be added
# to the flags, overriding this -O0.
if {[lsearch -exact $options hip] != -1} {
lappend new_options "early_flags=-O0"
}
# Because we link with libraries using their basename, we may need
# (depending on the platform) to set a special rpath value, to allow
# the executable to find the libraries it depends on.

94
gdb/testsuite/lib/rocm.exp Normal file
View file

@ -0,0 +1,94 @@
# Copyright (C) 2019-2022 Free Software Foundation, Inc.
#
# 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/>.
#
# Support library for testing ROCm (AMD GPU) GDB features.
proc skip_hipcc_tests { } {
# Only the native target supports ROCm debugging. E.g., when
# testing against GDBserver, there's no point in running the ROCm
# tests.
if {[target_info gdb_protocol] != ""} {
return 1
}
return 0
}
# 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
# Acquire lock file LOCKFILE. Tries forever until the lock file is
# successfully created.
proc lock_file_acquire {lockfile} {
verbose -log "acquiring lock file: $::subdir/${::gdb_test_file_name}.exp"
while {true} {
if {![catch {open $lockfile {WRONLY CREAT EXCL}} rc]} {
set msg "locked by $::subdir/${::gdb_test_file_name}.exp"
verbose -log "lock file: $msg"
# For debugging, put info in the lockfile about who owns
# it.
puts $rc $msg
flush $rc
return [list $rc $lockfile]
}
after 10
}
}
# Release a lock file.
proc lock_file_release {info} {
verbose -log "releasing lock file: $::subdir/${::gdb_test_file_name}.exp"
if {![catch {fconfigure [lindex $info 0]}]} {
if {![catch {
close [lindex $info 0]
file delete -force [lindex $info 1]
} rc]} {
return ""
} else {
return -code error "Error releasing lockfile: '$rc'"
}
} else {
error "invalid lock"
}
}
# Run body under the GPU lock. Also calls gdb_exit before releasing
# the GPU lock.
proc with_rocm_gpu_lock { body } {
if {[info exists ::GDB_PARALLEL]} {
set lock_rc [lock_file_acquire $::gpu_lock_filename]
}
set code [catch {uplevel 1 $body} result]
# In case BODY returned early due to some testcase failing, and
# left GDB running, debugging the GPU.
gdb_exit
if {[info exists ::GDB_PARALLEL]} {
lock_file_release $lock_rc
}
if {$code == 1} {
global errorInfo errorCode
return -code $code -errorinfo $errorInfo -errorcode $errorCode $result
} else {
return -code $code $result
}
}