public inbox for gdb-cvs@sourceware.org
help / color / mirror / Atom feed
* [binutils-gdb] gdb: initial support for ROCm platform (AMDGPU) debugging
@ 2023-02-02 15:09 Simon Marchi
0 siblings, 0 replies; only message in thread
From: Simon Marchi @ 2023-02-02 15:09 UTC (permalink / raw)
To: gdb-cvs
https://sourceware.org/git/gitweb.cgi?p=binutils-gdb.git;h=18b4d0736bc570c6d2e3e5f6ebc2ad4617d93847
commit 18b4d0736bc570c6d2e3e5f6ebc2ad4617d93847
Author: Simon Marchi <simon.marchi@polymtl.ca>
Date: Tue Jan 3 15:07:07 2023 -0500
gdb: initial support for ROCm platform (AMDGPU) debugging
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>
Diff:
---
gdb/Makefile.in | 17 +-
gdb/NEWS | 7 +
gdb/README | 15 +
gdb/amd-dbgapi-target.c | 1966 +++++++++++++++++++++++++++++++++++++
gdb/amd-dbgapi-target.h | 116 +++
gdb/amdgpu-tdep.c | 1367 ++++++++++++++++++++++++++
gdb/amdgpu-tdep.h | 93 ++
gdb/configure | 425 +++++---
gdb/configure.ac | 52 +
gdb/configure.tgt | 23 +-
gdb/doc/gdb.texinfo | 291 ++++++
gdb/regcache.c | 3 +-
gdb/solib-rocm.c | 679 +++++++++++++
gdb/testsuite/gdb.rocm/simple.cpp | 48 +
gdb/testsuite/gdb.rocm/simple.exp | 52 +
gdb/testsuite/lib/future.exp | 38 +
gdb/testsuite/lib/gdb.exp | 7 +
gdb/testsuite/lib/rocm.exp | 94 ++
18 files changed, 5155 insertions(+), 138 deletions(-)
diff --git a/gdb/Makefile.in b/gdb/Makefile.in
index c3711a0d2f5..049a14fe40a 100644
--- a/gdb/Makefile.in
+++ b/gdb/Makefile.in
@@ -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 \
diff --git a/gdb/NEWS b/gdb/NEWS
index 445d28efed9..882ea4cda36 100644
--- a/gdb/NEWS
+++ b/gdb/NEWS
@@ -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
diff --git a/gdb/README b/gdb/README
index fbe480f0d60..9699f4890c6 100644
--- a/gdb/README
+++ b/gdb/README
@@ -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
diff --git a/gdb/amd-dbgapi-target.c b/gdb/amd-dbgapi-target.c
new file mode 100644
index 00000000000..5f7de52a1a5
--- /dev/null
+++ b/gdb/amd-dbgapi-target.c
@@ -0,0 +1,1966 @@
+/* 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/>. */
+
+#include "defs.h"
+
+#include "amd-dbgapi-target.h"
+#include "amdgpu-tdep.h"
+#include "async-event.h"
+#include "cli/cli-cmds.h"
+#include "cli/cli-style.h"
+#include "inf-loop.h"
+#include "inferior.h"
+#include "objfiles.h"
+#include "observable.h"
+#include "registry.h"
+#include "solib.h"
+#include "target.h"
+
+/* When true, print debug messages relating to the amd-dbgapi target. */
+
+static bool debug_amd_dbgapi = false;
+
+/* Make a copy of S styled in green. */
+
+static std::string
+make_green (const char *s)
+{
+ cli_style_option style (nullptr, ui_file_style::GREEN);
+ string_file sf (true);
+ gdb_printf (&sf, "%ps", styled_string (style.style(), s));
+ return sf.release ();
+}
+
+/* Debug module names. "amd-dbgapi" is for the target debug messages (this
+ file), whereas "amd-dbgapi-lib" is for logging messages output by the
+ amd-dbgapi library. */
+
+static const char *amd_dbgapi_debug_module_unstyled = "amd-dbgapi";
+static const char *amd_dbgapi_lib_debug_module_unstyled
+ = "amd-dbgapi-lib";
+
+/* Styled variants of the above. */
+
+static const std::string amd_dbgapi_debug_module_styled
+ = make_green (amd_dbgapi_debug_module_unstyled);
+static const std::string amd_dbgapi_lib_debug_module_styled
+ = make_green (amd_dbgapi_lib_debug_module_unstyled);
+
+/* Return the styled or unstyled variant of the amd-dbgapi module name,
+ depending on whether gdb_stdlog can emit colors. */
+
+static const char *
+amd_dbgapi_debug_module ()
+{
+ if (gdb_stdlog->can_emit_style_escape ())
+ return amd_dbgapi_debug_module_styled.c_str ();
+ else
+ return amd_dbgapi_debug_module_unstyled;
+}
+
+/* Same as the above, but for the amd-dbgapi-lib module name. */
+
+static const char *
+amd_dbgapi_lib_debug_module ()
+{
+ if (gdb_stdlog->can_emit_style_escape ())
+ return amd_dbgapi_lib_debug_module_styled.c_str ();
+ else
+ return amd_dbgapi_lib_debug_module_unstyled;
+}
+
+/* Print an amd-dbgapi debug statement. */
+
+#define amd_dbgapi_debug_printf(fmt, ...) \
+ debug_prefixed_printf_cond (debug_amd_dbgapi, \
+ amd_dbgapi_debug_module (), \
+ fmt, ##__VA_ARGS__)
+
+/* Print amd-dbgapi start/end debug statements. */
+
+#define AMD_DBGAPI_SCOPED_DEBUG_START_END(fmt, ...) \
+ scoped_debug_start_end (debug_infrun, amd_dbgapi_debug_module (), \
+ fmt, ##__VA_ARGS__)
+
+/* inferior_created observer token. */
+
+static gdb::observers::token amd_dbgapi_target_inferior_created_observer_token;
+
+const gdb::observers::token &
+get_amd_dbgapi_target_inferior_created_observer_token ()
+{
+ return amd_dbgapi_target_inferior_created_observer_token;
+}
+
+
+/* Big enough to hold the size of the largest register in bytes. */
+#define AMDGPU_MAX_REGISTER_SIZE 256
+
+/* amd-dbgapi-specific inferior data. */
+
+struct amd_dbgapi_inferior_info
+{
+ explicit amd_dbgapi_inferior_info (inferior *inf)
+ : inf (inf)
+ {}
+
+ /* Backlink to inferior. */
+ inferior *inf;
+
+ /* The amd_dbgapi_process_id for this inferior. */
+ amd_dbgapi_process_id_t process_id = AMD_DBGAPI_PROCESS_NONE;
+
+ /* The amd_dbgapi_notifier_t for this inferior. */
+ amd_dbgapi_notifier_t notifier = -1;
+
+ /* The status of the inferior's runtime support. */
+ amd_dbgapi_runtime_state_t runtime_state = AMD_DBGAPI_RUNTIME_STATE_UNLOADED;
+
+ /* This value mirrors the current "forward progress needed" value for this
+ process in amd-dbgapi. It is used to avoid unnecessary calls to
+ amd_dbgapi_process_set_progress, to reduce the noise in the logs.
+
+ Initialized to true, since that's the default in amd-dbgapi too. */
+ bool forward_progress_required = true;
+
+ std::unordered_map<decltype (amd_dbgapi_breakpoint_id_t::handle),
+ struct breakpoint *>
+ breakpoint_map;
+
+ /* List of pending events the amd-dbgapi target retrieved from the dbgapi. */
+ std::list<std::pair<ptid_t, target_waitstatus>> wave_events;
+};
+
+static amd_dbgapi_event_id_t process_event_queue
+ (amd_dbgapi_process_id_t process_id = AMD_DBGAPI_PROCESS_NONE,
+ amd_dbgapi_event_kind_t until_event_kind = AMD_DBGAPI_EVENT_KIND_NONE);
+
+static const target_info amd_dbgapi_target_info = {
+ "amd-dbgapi",
+ N_("AMD Debugger API"),
+ N_("GPU debugging using the AMD Debugger API")
+};
+
+static amd_dbgapi_log_level_t get_debug_amd_dbgapi_lib_log_level ();
+
+struct amd_dbgapi_target final : public target_ops
+{
+ const target_info &
+ info () const override
+ {
+ return amd_dbgapi_target_info;
+ }
+ strata
+ stratum () const override
+ {
+ return arch_stratum;
+ }
+
+ void close () override;
+ void mourn_inferior () override;
+ void detach (inferior *inf, int from_tty) override;
+
+ void async (bool enable) override;
+
+ bool has_pending_events () override;
+ ptid_t wait (ptid_t, struct target_waitstatus *, target_wait_flags) override;
+ void resume (ptid_t, int, enum gdb_signal) override;
+ void commit_resumed () override;
+ void stop (ptid_t ptid) override;
+
+ void fetch_registers (struct regcache *, int) override;
+ void store_registers (struct regcache *, int) override;
+
+ void update_thread_list () override;
+
+ struct gdbarch *thread_architecture (ptid_t) override;
+
+ void thread_events (int enable) override;
+
+ std::string pid_to_str (ptid_t ptid) override;
+
+ const char *thread_name (thread_info *tp) override;
+
+ const char *extra_thread_info (thread_info *tp) override;
+
+ bool thread_alive (ptid_t ptid) override;
+
+ enum target_xfer_status xfer_partial (enum target_object object,
+ const char *annex, gdb_byte *readbuf,
+ const gdb_byte *writebuf,
+ ULONGEST offset, ULONGEST len,
+ ULONGEST *xfered_len) override;
+
+ bool stopped_by_watchpoint () override;
+
+ bool stopped_by_sw_breakpoint () override;
+ bool stopped_by_hw_breakpoint () override;
+
+private:
+ /* True if we must report thread events. */
+ bool m_report_thread_events = false;
+
+ /* Cache for the last value returned by thread_architecture. */
+ gdbarch *m_cached_arch = nullptr;
+ ptid_t::tid_type m_cached_arch_tid = 0;
+};
+
+static struct amd_dbgapi_target the_amd_dbgapi_target;
+
+/* Per-inferior data key. */
+
+static const registry<inferior>::key<amd_dbgapi_inferior_info>
+ amd_dbgapi_inferior_data;
+
+/* The async event handler registered with the event loop, indicating that we
+ might have events to report to the core and that we'd like our wait method
+ to be called.
+
+ This is nullptr when async is disabled and non-nullptr when async is
+ enabled.
+
+ It is marked when a notifier fd tells us there's an event available. The
+ callback triggers handle_inferior_event in order to pull the event from
+ amd-dbgapi and handle it. */
+
+static async_event_handler *amd_dbgapi_async_event_handler = nullptr;
+
+/* Return the target id string for a given wave. */
+
+static std::string
+wave_target_id_string (amd_dbgapi_wave_id_t wave_id)
+{
+ amd_dbgapi_dispatch_id_t dispatch_id;
+ amd_dbgapi_queue_id_t queue_id;
+ amd_dbgapi_agent_id_t agent_id;
+ uint32_t group_ids[3], wave_in_group;
+ std::string str = "AMDGPU Wave";
+
+ amd_dbgapi_status_t status
+ = amd_dbgapi_wave_get_info (wave_id, AMD_DBGAPI_WAVE_INFO_AGENT,
+ sizeof (agent_id), &agent_id);
+ str += (status == AMD_DBGAPI_STATUS_SUCCESS
+ ? string_printf (" %ld", agent_id.handle)
+ : " ?");
+
+ status = amd_dbgapi_wave_get_info (wave_id, AMD_DBGAPI_WAVE_INFO_QUEUE,
+ sizeof (queue_id), &queue_id);
+ str += (status == AMD_DBGAPI_STATUS_SUCCESS
+ ? string_printf (":%ld", queue_id.handle)
+ : ":?");
+
+ status = amd_dbgapi_wave_get_info (wave_id, AMD_DBGAPI_WAVE_INFO_DISPATCH,
+ sizeof (dispatch_id), &dispatch_id);
+ str += (status == AMD_DBGAPI_STATUS_SUCCESS
+ ? string_printf (":%ld", dispatch_id.handle)
+ : ":?");
+
+ str += string_printf (":%ld", wave_id.handle);
+
+ status = amd_dbgapi_wave_get_info (wave_id,
+ AMD_DBGAPI_WAVE_INFO_WORKGROUP_COORD,
+ sizeof (group_ids), &group_ids);
+ str += (status == AMD_DBGAPI_STATUS_SUCCESS
+ ? string_printf (" (%d,%d,%d)", group_ids[0], group_ids[1],
+ group_ids[2])
+ : " (?,?,?)");
+
+ status = amd_dbgapi_wave_get_info
+ (wave_id, AMD_DBGAPI_WAVE_INFO_WAVE_NUMBER_IN_WORKGROUP,
+ sizeof (wave_in_group), &wave_in_group);
+ str += (status == AMD_DBGAPI_STATUS_SUCCESS
+ ? string_printf ("/%d", wave_in_group)
+ : "/?");
+
+ return str;
+}
+
+/* Clear our async event handler. */
+
+static void
+async_event_handler_clear ()
+{
+ gdb_assert (amd_dbgapi_async_event_handler != nullptr);
+ clear_async_event_handler (amd_dbgapi_async_event_handler);
+}
+
+/* Mark our async event handler. */
+
+static void
+async_event_handler_mark ()
+{
+ gdb_assert (amd_dbgapi_async_event_handler != nullptr);
+ mark_async_event_handler (amd_dbgapi_async_event_handler);
+}
+
+/* Fetch the amd_dbgapi_inferior_info data for the given inferior. */
+
+static struct amd_dbgapi_inferior_info *
+get_amd_dbgapi_inferior_info (struct inferior *inferior)
+{
+ amd_dbgapi_inferior_info *info = amd_dbgapi_inferior_data.get (inferior);
+
+ if (info == nullptr)
+ info = amd_dbgapi_inferior_data.emplace (inferior, inferior);
+
+ return info;
+}
+
+/* Set forward progress requirement to REQUIRE for all processes of PROC_TARGET
+ matching PTID. */
+
+static void
+require_forward_progress (ptid_t ptid, process_stratum_target *proc_target,
+ bool require)
+{
+ for (inferior *inf : all_inferiors (proc_target))
+ {
+ if (ptid != minus_one_ptid && inf->pid != ptid.pid ())
+ continue;
+
+ amd_dbgapi_inferior_info *info = get_amd_dbgapi_inferior_info (inf);
+
+ if (info->process_id == AMD_DBGAPI_PROCESS_NONE)
+ continue;
+
+ /* Don't do unnecessary calls to amd-dbgapi to avoid polluting the logs. */
+ if (info->forward_progress_required == require)
+ continue;
+
+ amd_dbgapi_status_t status
+ = amd_dbgapi_process_set_progress
+ (info->process_id, (require
+ ? AMD_DBGAPI_PROGRESS_NORMAL
+ : AMD_DBGAPI_PROGRESS_NO_FORWARD));
+ gdb_assert (status == AMD_DBGAPI_STATUS_SUCCESS);
+
+ info->forward_progress_required = require;
+
+ /* If ptid targets a single inferior and we have found it, no need to
+ continue. */
+ if (ptid != minus_one_ptid)
+ break;
+ }
+}
+
+/* See amd-dbgapi-target.h. */
+
+amd_dbgapi_process_id_t
+get_amd_dbgapi_process_id (inferior *inf)
+{
+ return get_amd_dbgapi_inferior_info (inf)->process_id;
+}
+
+/* A breakpoint dbgapi wants us to insert, to handle shared library
+ loading/unloading. */
+
+struct amd_dbgapi_target_breakpoint : public code_breakpoint
+{
+ amd_dbgapi_target_breakpoint (struct gdbarch *gdbarch, CORE_ADDR address)
+ : code_breakpoint (gdbarch, bp_breakpoint)
+ {
+ symtab_and_line sal;
+ sal.pc = address;
+ sal.section = find_pc_overlay (sal.pc);
+ sal.pspace = current_program_space;
+ add_location (sal);
+
+ pspace = current_program_space;
+ disposition = disp_donttouch;
+ }
+
+ void re_set () override;
+ void check_status (struct bpstat *bs) override;
+};
+
+void
+amd_dbgapi_target_breakpoint::re_set ()
+{
+ /* Nothing. */
+}
+
+void
+amd_dbgapi_target_breakpoint::check_status (struct bpstat *bs)
+{
+ inferior *inf = current_inferior ();
+ amd_dbgapi_inferior_info *info = get_amd_dbgapi_inferior_info (inf);
+ amd_dbgapi_status_t status;
+
+ bs->stop = 0;
+ bs->print_it = print_it_noop;
+
+ /* Find the address the breakpoint is set at. */
+ auto match_breakpoint
+ = [bs] (const decltype (info->breakpoint_map)::value_type &value)
+ { return value.second == bs->breakpoint_at; };
+ auto it
+ = std::find_if (info->breakpoint_map.begin (), info->breakpoint_map.end (),
+ match_breakpoint);
+
+ if (it == info->breakpoint_map.end ())
+ error (_("Could not find breakpoint_id for breakpoint at %s"),
+ paddress (inf->gdbarch, bs->bp_location_at->address));
+
+ amd_dbgapi_breakpoint_id_t breakpoint_id { it->first };
+ amd_dbgapi_breakpoint_action_t action;
+
+ status = amd_dbgapi_report_breakpoint_hit
+ (breakpoint_id,
+ reinterpret_cast<amd_dbgapi_client_thread_id_t> (inferior_thread ()),
+ &action);
+
+ if (status != AMD_DBGAPI_STATUS_SUCCESS)
+ error (_("amd_dbgapi_report_breakpoint_hit failed for breakpoint %ld "
+ "at %s (%s)"),
+ breakpoint_id.handle, paddress (inf->gdbarch, bs->bp_location_at->address),
+ get_status_string (status));
+
+ if (action == AMD_DBGAPI_BREAKPOINT_ACTION_RESUME)
+ return;
+
+ /* If the action is AMD_DBGAPI_BREAKPOINT_ACTION_HALT, we need to wait until
+ a breakpoint resume event for this breakpoint_id is seen. */
+ amd_dbgapi_event_id_t resume_event_id
+ = process_event_queue (info->process_id,
+ AMD_DBGAPI_EVENT_KIND_BREAKPOINT_RESUME);
+
+ /* We should always get a breakpoint_resume event after processing all
+ events generated by reporting the breakpoint hit. */
+ gdb_assert (resume_event_id != AMD_DBGAPI_EVENT_NONE);
+
+ amd_dbgapi_breakpoint_id_t resume_breakpoint_id;
+ status = amd_dbgapi_event_get_info (resume_event_id,
+ AMD_DBGAPI_EVENT_INFO_BREAKPOINT,
+ sizeof (resume_breakpoint_id),
+ &resume_breakpoint_id);
+
+ if (status != AMD_DBGAPI_STATUS_SUCCESS)
+ error (_("amd_dbgapi_event_get_info failed (%s)"), get_status_string (status));
+
+ /* The debugger API guarantees that [breakpoint_hit...resume_breakpoint]
+ sequences cannot interleave, so this breakpoint resume event must be
+ for our breakpoint_id. */
+ if (resume_breakpoint_id != breakpoint_id)
+ error (_("breakpoint resume event is not for this breakpoint. "
+ "Expected breakpoint_%ld, got breakpoint_%ld"),
+ breakpoint_id.handle, resume_breakpoint_id.handle);
+
+ amd_dbgapi_event_processed (resume_event_id);
+}
+
+bool
+amd_dbgapi_target::thread_alive (ptid_t ptid)
+{
+ if (!ptid_is_gpu (ptid))
+ return beneath ()->thread_alive (ptid);
+
+ /* Check that the wave_id is valid. */
+
+ amd_dbgapi_wave_state_t state;
+ amd_dbgapi_status_t status
+ = amd_dbgapi_wave_get_info (get_amd_dbgapi_wave_id (ptid),
+ AMD_DBGAPI_WAVE_INFO_STATE, sizeof (state),
+ &state);
+ return status == AMD_DBGAPI_STATUS_SUCCESS;
+}
+
+const char *
+amd_dbgapi_target::thread_name (thread_info *tp)
+{
+ if (!ptid_is_gpu (tp->ptid))
+ return beneath ()->thread_name (tp);
+
+ return nullptr;
+}
+
+std::string
+amd_dbgapi_target::pid_to_str (ptid_t ptid)
+{
+ if (!ptid_is_gpu (ptid))
+ return beneath ()->pid_to_str (ptid);
+
+ return wave_target_id_string (get_amd_dbgapi_wave_id (ptid));
+}
+
+const char *
+amd_dbgapi_target::extra_thread_info (thread_info *tp)
+{
+ if (!ptid_is_gpu (tp->ptid))
+ beneath ()->extra_thread_info (tp);
+
+ return nullptr;
+}
+
+target_xfer_status
+amd_dbgapi_target::xfer_partial (enum target_object object, const char *annex,
+ gdb_byte *readbuf, const gdb_byte *writebuf,
+ ULONGEST offset, ULONGEST requested_len,
+ ULONGEST *xfered_len)
+{
+ gdb::optional<scoped_restore_current_thread> maybe_restore_thread;
+
+ if (!ptid_is_gpu (inferior_ptid))
+ return beneath ()->xfer_partial (object, annex, readbuf, writebuf, offset,
+ requested_len, xfered_len);
+
+ gdb_assert (requested_len > 0);
+ gdb_assert (xfered_len != nullptr);
+
+ if (object != TARGET_OBJECT_MEMORY)
+ return TARGET_XFER_E_IO;
+
+ amd_dbgapi_process_id_t process_id
+ = get_amd_dbgapi_process_id (current_inferior ());
+ amd_dbgapi_wave_id_t wave_id = get_amd_dbgapi_wave_id (inferior_ptid);
+
+ size_t len = requested_len;
+ amd_dbgapi_status_t status;
+
+ if (readbuf != nullptr)
+ status = amd_dbgapi_read_memory (process_id, wave_id, 0,
+ AMD_DBGAPI_ADDRESS_SPACE_GLOBAL,
+ offset, &len, readbuf);
+ else
+ status = amd_dbgapi_write_memory (process_id, wave_id, 0,
+ AMD_DBGAPI_ADDRESS_SPACE_GLOBAL,
+ offset, &len, writebuf);
+
+ if (status != AMD_DBGAPI_STATUS_SUCCESS)
+ return TARGET_XFER_E_IO;
+
+ *xfered_len = len;
+ return TARGET_XFER_OK;
+}
+
+bool
+amd_dbgapi_target::stopped_by_watchpoint ()
+{
+ if (!ptid_is_gpu (inferior_ptid))
+ return beneath ()->stopped_by_watchpoint ();
+
+ return false;
+}
+
+void
+amd_dbgapi_target::resume (ptid_t scope_ptid, int step, enum gdb_signal signo)
+{
+ amd_dbgapi_debug_printf ("scope_ptid = %s", scope_ptid.to_string ().c_str ());
+
+ /* The amd_dbgapi_exceptions_t matching SIGNO will only be used if the
+ thread which is the target of the signal SIGNO is a GPU thread. If so,
+ make sure that there is a corresponding amd_dbgapi_exceptions_t for SIGNO
+ before we try to resume any thread. */
+ amd_dbgapi_exceptions_t exception = AMD_DBGAPI_EXCEPTION_NONE;
+ if (ptid_is_gpu (inferior_ptid))
+ {
+ switch (signo)
+ {
+ case GDB_SIGNAL_BUS:
+ exception = AMD_DBGAPI_EXCEPTION_WAVE_APERTURE_VIOLATION;
+ break;
+ case GDB_SIGNAL_SEGV:
+ exception = AMD_DBGAPI_EXCEPTION_WAVE_MEMORY_VIOLATION;
+ break;
+ case GDB_SIGNAL_ILL:
+ exception = AMD_DBGAPI_EXCEPTION_WAVE_ILLEGAL_INSTRUCTION;
+ break;
+ case GDB_SIGNAL_FPE:
+ exception = AMD_DBGAPI_EXCEPTION_WAVE_MATH_ERROR;
+ break;
+ case GDB_SIGNAL_ABRT:
+ exception = AMD_DBGAPI_EXCEPTION_WAVE_ABORT;
+ break;
+ case GDB_SIGNAL_TRAP:
+ exception = AMD_DBGAPI_EXCEPTION_WAVE_TRAP;
+ break;
+ case GDB_SIGNAL_0:
+ exception = AMD_DBGAPI_EXCEPTION_NONE;
+ break;
+ default:
+ error (_("Resuming with signal %s is not supported by this agent."),
+ gdb_signal_to_name (signo));
+ }
+ }
+
+ if (!ptid_is_gpu (inferior_ptid) || scope_ptid != inferior_ptid)
+ {
+ beneath ()->resume (scope_ptid, step, signo);
+
+ /* If the request is for a single thread, we are done. */
+ if (scope_ptid == inferior_ptid)
+ return;
+ }
+
+ process_stratum_target *proc_target = current_inferior ()->process_target ();
+
+ /* Disable forward progress requirement. */
+ require_forward_progress (scope_ptid, proc_target, false);
+
+ for (thread_info *thread : all_non_exited_threads (proc_target, scope_ptid))
+ {
+ if (!ptid_is_gpu (thread->ptid))
+ continue;
+
+ amd_dbgapi_wave_id_t wave_id = get_amd_dbgapi_wave_id (thread->ptid);
+ amd_dbgapi_status_t status;
+ if (thread->ptid == inferior_ptid)
+ status = amd_dbgapi_wave_resume (wave_id,
+ (step
+ ? AMD_DBGAPI_RESUME_MODE_SINGLE_STEP
+ : AMD_DBGAPI_RESUME_MODE_NORMAL),
+ exception);
+ else
+ status = amd_dbgapi_wave_resume (wave_id, AMD_DBGAPI_RESUME_MODE_NORMAL,
+ AMD_DBGAPI_EXCEPTION_NONE);
+
+ if (status != AMD_DBGAPI_STATUS_SUCCESS
+ /* Ignore the error that wave is no longer valid as that could
+ indicate that the process has exited. GDB treats resuming a
+ thread that no longer exists as being successful. */
+ && status != AMD_DBGAPI_STATUS_ERROR_INVALID_WAVE_ID)
+ error (_("wave_resume for wave_%ld failed (%s)"), wave_id.handle,
+ get_status_string (status));
+ }
+}
+
+void
+amd_dbgapi_target::commit_resumed ()
+{
+ amd_dbgapi_debug_printf ("called");
+
+ beneath ()->commit_resumed ();
+
+ process_stratum_target *proc_target = current_inferior ()->process_target ();
+ require_forward_progress (minus_one_ptid, proc_target, true);
+}
+
+void
+amd_dbgapi_target::stop (ptid_t ptid)
+{
+ amd_dbgapi_debug_printf ("ptid = %s", ptid.to_string ().c_str ());
+
+ bool many_threads = ptid == minus_one_ptid || ptid.is_pid ();
+
+ if (!ptid_is_gpu (ptid) || many_threads)
+ {
+ beneath ()->stop (ptid);
+
+ /* The request is for a single thread, we are done. */
+ if (!many_threads)
+ return;
+ }
+
+ auto stop_one_thread = [this] (thread_info *thread)
+ {
+ gdb_assert (thread != nullptr);
+
+ amd_dbgapi_wave_id_t wave_id = get_amd_dbgapi_wave_id (thread->ptid);
+ amd_dbgapi_wave_state_t state;
+ amd_dbgapi_status_t status
+ = amd_dbgapi_wave_get_info (wave_id, AMD_DBGAPI_WAVE_INFO_STATE,
+ sizeof (state), &state);
+ if (status == AMD_DBGAPI_STATUS_SUCCESS)
+ {
+ /* If the wave is already known to be stopped then do nothing. */
+ if (state == AMD_DBGAPI_WAVE_STATE_STOP)
+ return;
+
+ status = amd_dbgapi_wave_stop (wave_id);
+ if (status == AMD_DBGAPI_STATUS_SUCCESS)
+ return;
+
+ if (status != AMD_DBGAPI_STATUS_ERROR_INVALID_WAVE_ID)
+ error (_("wave_stop for wave_%ld failed (%s)"), wave_id.handle,
+ get_status_string (status));
+ }
+ else if (status != AMD_DBGAPI_STATUS_ERROR_INVALID_WAVE_ID)
+ error (_("wave_get_info for wave_%ld failed (%s)"), wave_id.handle,
+ get_status_string (status));
+
+ /* The status is AMD_DBGAPI_STATUS_ERROR_INVALID_WAVE_ID. The wave
+ could have terminated since the last time the wave list was
+ refreshed. */
+
+ if (m_report_thread_events)
+ {
+ get_amd_dbgapi_inferior_info (thread->inf)->wave_events.emplace_back
+ (thread->ptid, target_waitstatus ().set_thread_exited (0));
+
+ if (target_is_async_p ())
+ async_event_handler_mark ();
+ }
+
+ delete_thread_silent (thread);
+ };
+
+ process_stratum_target *proc_target = current_inferior ()->process_target ();
+
+ /* Disable forward progress requirement. */
+ require_forward_progress (ptid, proc_target, false);
+
+ if (!many_threads)
+ {
+ /* No need to iterate all non-exited threads if the request is to stop a
+ specific thread. */
+ stop_one_thread (find_thread_ptid (proc_target, ptid));
+ return;
+ }
+
+ for (auto *inf : all_inferiors (proc_target))
+ /* Use the threads_safe iterator since stop_one_thread may delete the
+ thread if it has exited. */
+ for (auto *thread : inf->threads_safe ())
+ if (thread->state != THREAD_EXITED && thread->ptid.matches (ptid)
+ && ptid_is_gpu (thread->ptid))
+ stop_one_thread (thread);
+}
+
+/* Callback for our async event handler. */
+
+static void
+handle_target_event (gdb_client_data client_data)
+{
+ inferior_event_handler (INF_REG_EVENT);
+}
+
+struct scoped_amd_dbgapi_event_processed
+{
+ scoped_amd_dbgapi_event_processed (amd_dbgapi_event_id_t event_id)
+ : m_event_id (event_id)
+ {
+ gdb_assert (event_id != AMD_DBGAPI_EVENT_NONE);
+ }
+
+ ~scoped_amd_dbgapi_event_processed ()
+ {
+ amd_dbgapi_status_t status = amd_dbgapi_event_processed (m_event_id);
+ if (status != AMD_DBGAPI_STATUS_SUCCESS)
+ warning (_("Failed to acknowledge amd-dbgapi event %" PRIu64),
+ m_event_id.handle);
+ }
+
+ DISABLE_COPY_AND_ASSIGN (scoped_amd_dbgapi_event_processed);
+
+private:
+ amd_dbgapi_event_id_t m_event_id;
+};
+
+/* Called when a dbgapi notifier fd is readable. CLIENT_DATA is the
+ amd_dbgapi_inferior_info object corresponding to the notifier. */
+
+static void
+dbgapi_notifier_handler (int err, gdb_client_data client_data)
+{
+ amd_dbgapi_inferior_info *info = (amd_dbgapi_inferior_info *) client_data;
+ int ret;
+
+ /* Drain the notifier pipe. */
+ do
+ {
+ char buf;
+ ret = read (info->notifier, &buf, 1);
+ }
+ while (ret >= 0 || (ret == -1 && errno == EINTR));
+
+ if (info->inf->target_is_pushed (&the_amd_dbgapi_target))
+ {
+ /* The amd-dbgapi target is pushed: signal our async handler, the event
+ will be consumed through our wait method. */
+
+ async_event_handler_mark ();
+ }
+ else
+ {
+ /* The amd-dbgapi target is not pushed: if there's an event, the only
+ expected one is one of the RUNTIME kind. If the event tells us the
+ inferior as activated the ROCm runtime, push the amd-dbgapi
+ target. */
+
+ amd_dbgapi_event_id_t event_id;
+ amd_dbgapi_event_kind_t event_kind;
+ amd_dbgapi_status_t status
+ = amd_dbgapi_process_next_pending_event (info->process_id, &event_id,
+ &event_kind);
+ if (status != AMD_DBGAPI_STATUS_SUCCESS)
+ error (_("next_pending_event failed (%s)"), get_status_string (status));
+
+ if (event_id == AMD_DBGAPI_EVENT_NONE)
+ return;
+
+ gdb_assert (event_kind == AMD_DBGAPI_EVENT_KIND_RUNTIME);
+
+ scoped_amd_dbgapi_event_processed mark_event_processed (event_id);
+
+ amd_dbgapi_runtime_state_t runtime_state;
+ status = amd_dbgapi_event_get_info (event_id,
+ AMD_DBGAPI_EVENT_INFO_RUNTIME_STATE,
+ sizeof (runtime_state),
+ &runtime_state);
+ if (status != AMD_DBGAPI_STATUS_SUCCESS)
+ error (_("event_get_info for event_%ld failed (%s)"),
+ event_id.handle, get_status_string (status));
+
+ switch (runtime_state)
+ {
+ case AMD_DBGAPI_RUNTIME_STATE_LOADED_SUCCESS:
+ gdb_assert (info->runtime_state == AMD_DBGAPI_RUNTIME_STATE_UNLOADED);
+ info->runtime_state = runtime_state;
+ amd_dbgapi_debug_printf ("pushing amd-dbgapi target");
+ info->inf->push_target (&the_amd_dbgapi_target);
+
+ /* The underlying target will already be async if we are running, but not if
+ we are attaching. */
+ if (info->inf->process_target ()->is_async_p ())
+ {
+ scoped_restore_current_thread restore_thread;
+ switch_to_inferior_no_thread (info->inf);
+
+ /* Make sure our async event handler is created. */
+ target_async (true);
+ }
+ break;
+
+ case AMD_DBGAPI_RUNTIME_STATE_UNLOADED:
+ gdb_assert (info->runtime_state
+ == AMD_DBGAPI_RUNTIME_STATE_LOADED_ERROR_RESTRICTION);
+ info->runtime_state = runtime_state;
+ break;
+
+ case AMD_DBGAPI_RUNTIME_STATE_LOADED_ERROR_RESTRICTION:
+ gdb_assert (info->runtime_state == AMD_DBGAPI_RUNTIME_STATE_UNLOADED);
+ info->runtime_state = runtime_state;
+ warning (_("amd-dbgapi: unable to enable GPU debugging "
+ "due to a restriction error"));
+ break;
+ }
+ }
+}
+
+void
+amd_dbgapi_target::async (bool enable)
+{
+ beneath ()->async (enable);
+
+ if (enable)
+ {
+ if (amd_dbgapi_async_event_handler != nullptr)
+ {
+ /* Already enabled. */
+ return;
+ }
+
+ /* The library gives us one notifier file descriptor per inferior (even
+ the ones that have not yet loaded their runtime). Register them
+ all with the event loop. */
+ process_stratum_target *proc_target
+ = current_inferior ()->process_target ();
+
+ for (inferior *inf : all_non_exited_inferiors (proc_target))
+ {
+ amd_dbgapi_inferior_info *info = get_amd_dbgapi_inferior_info (inf);
+
+ if (info->notifier != -1)
+ add_file_handler (info->notifier, dbgapi_notifier_handler, info,
+ string_printf ("amd-dbgapi notifier for pid %d",
+ inf->pid));
+ }
+
+ amd_dbgapi_async_event_handler
+ = create_async_event_handler (handle_target_event, nullptr,
+ "amd-dbgapi");
+
+ /* There may be pending events to handle. Tell the event loop to poll
+ them. */
+ async_event_handler_mark ();
+ }
+ else
+ {
+ if (amd_dbgapi_async_event_handler == nullptr)
+ return;
+
+ for (inferior *inf : all_inferiors ())
+ {
+ amd_dbgapi_inferior_info *info = get_amd_dbgapi_inferior_info (inf);
+
+ if (info->notifier != -1)
+ delete_file_handler (info->notifier);
+ }
+
+ delete_async_event_handler (&amd_dbgapi_async_event_handler);
+ }
+}
+
+/* Make a ptid for a GPU wave. See comment on ptid_is_gpu for more details. */
+
+static ptid_t
+make_gpu_ptid (ptid_t::pid_type pid, amd_dbgapi_wave_id_t wave_id)
+{
+ return ptid_t (pid, 1, wave_id.handle);
+}
+
+/* Process an event that was just pulled out of the amd-dbgapi library. */
+
+static void
+process_one_event (amd_dbgapi_event_id_t event_id,
+ amd_dbgapi_event_kind_t event_kind)
+{
+ /* Automatically mark this event processed when going out of scope. */
+ scoped_amd_dbgapi_event_processed mark_event_processed (event_id);
+
+ amd_dbgapi_process_id_t process_id;
+ amd_dbgapi_status_t status
+ = amd_dbgapi_event_get_info (event_id, AMD_DBGAPI_EVENT_INFO_PROCESS,
+ sizeof (process_id), &process_id);
+ if (status != AMD_DBGAPI_STATUS_SUCCESS)
+ error (_("event_get_info for event_%ld failed (%s)"), event_id.handle,
+ get_status_string (status));
+
+ amd_dbgapi_os_process_id_t pid;
+ status = amd_dbgapi_process_get_info (process_id,
+ AMD_DBGAPI_PROCESS_INFO_OS_ID,
+ sizeof (pid), &pid);
+ if (status != AMD_DBGAPI_STATUS_SUCCESS)
+ error (_("process_get_info for process_%ld failed (%s)"),
+ process_id.handle, get_status_string (status));
+
+ auto *proc_target = current_inferior ()->process_target ();
+ inferior *inf = find_inferior_pid (proc_target, pid);
+ gdb_assert (inf != nullptr);
+ amd_dbgapi_inferior_info *info = get_amd_dbgapi_inferior_info (inf);
+
+ switch (event_kind)
+ {
+ case AMD_DBGAPI_EVENT_KIND_WAVE_COMMAND_TERMINATED:
+ case AMD_DBGAPI_EVENT_KIND_WAVE_STOP:
+ {
+ amd_dbgapi_wave_id_t wave_id;
+ status
+ = amd_dbgapi_event_get_info (event_id, AMD_DBGAPI_EVENT_INFO_WAVE,
+ sizeof (wave_id), &wave_id);
+ if (status != AMD_DBGAPI_STATUS_SUCCESS)
+ error (_("event_get_info for event_%ld failed (%s)"),
+ event_id.handle, get_status_string (status));
+
+ ptid_t event_ptid = make_gpu_ptid (pid, wave_id);
+ target_waitstatus ws;
+
+ amd_dbgapi_wave_stop_reasons_t stop_reason;
+ status = amd_dbgapi_wave_get_info (wave_id,
+ AMD_DBGAPI_WAVE_INFO_STOP_REASON,
+ sizeof (stop_reason), &stop_reason);
+ if (status == AMD_DBGAPI_STATUS_ERROR_INVALID_WAVE_ID
+ && event_kind == AMD_DBGAPI_EVENT_KIND_WAVE_COMMAND_TERMINATED)
+ ws.set_thread_exited (0);
+ else if (status == AMD_DBGAPI_STATUS_SUCCESS)
+ {
+ if (stop_reason & AMD_DBGAPI_WAVE_STOP_REASON_APERTURE_VIOLATION)
+ ws.set_stopped (GDB_SIGNAL_BUS);
+ else if (stop_reason
+ & AMD_DBGAPI_WAVE_STOP_REASON_MEMORY_VIOLATION)
+ ws.set_stopped (GDB_SIGNAL_SEGV);
+ else if (stop_reason
+ & AMD_DBGAPI_WAVE_STOP_REASON_ILLEGAL_INSTRUCTION)
+ ws.set_stopped (GDB_SIGNAL_ILL);
+ else if (stop_reason
+ & (AMD_DBGAPI_WAVE_STOP_REASON_FP_INPUT_DENORMAL
+ | AMD_DBGAPI_WAVE_STOP_REASON_FP_DIVIDE_BY_0
+ | AMD_DBGAPI_WAVE_STOP_REASON_FP_OVERFLOW
+ | AMD_DBGAPI_WAVE_STOP_REASON_FP_UNDERFLOW
+ | AMD_DBGAPI_WAVE_STOP_REASON_FP_INEXACT
+ | AMD_DBGAPI_WAVE_STOP_REASON_FP_INVALID_OPERATION
+ | AMD_DBGAPI_WAVE_STOP_REASON_INT_DIVIDE_BY_0))
+ ws.set_stopped (GDB_SIGNAL_FPE);
+ else if (stop_reason
+ & (AMD_DBGAPI_WAVE_STOP_REASON_BREAKPOINT
+ | AMD_DBGAPI_WAVE_STOP_REASON_WATCHPOINT
+ | AMD_DBGAPI_WAVE_STOP_REASON_SINGLE_STEP
+ | AMD_DBGAPI_WAVE_STOP_REASON_DEBUG_TRAP
+ | AMD_DBGAPI_WAVE_STOP_REASON_TRAP))
+ ws.set_stopped (GDB_SIGNAL_TRAP);
+ else if (stop_reason & AMD_DBGAPI_WAVE_STOP_REASON_ASSERT_TRAP)
+ ws.set_stopped (GDB_SIGNAL_ABRT);
+ else
+ ws.set_stopped (GDB_SIGNAL_0);
+
+ thread_info *thread = find_thread_ptid (proc_target, event_ptid);
+ if (thread == nullptr)
+ {
+ /* Silently create new GPU threads to avoid spamming the
+ terminal with thousands of "[New Thread ...]" messages. */
+ thread = add_thread_silent (proc_target, event_ptid);
+ set_running (proc_target, event_ptid, true);
+ set_executing (proc_target, event_ptid, true);
+ }
+
+ /* If the wave is stopped because of a software breakpoint, the
+ program counter needs to be adjusted so that it points to the
+ breakpoint instruction. */
+ if ((stop_reason & AMD_DBGAPI_WAVE_STOP_REASON_BREAKPOINT) != 0)
+ {
+ regcache *regcache = get_thread_regcache (thread);
+ gdbarch *gdbarch = regcache->arch ();
+
+ CORE_ADDR pc = regcache_read_pc (regcache);
+ CORE_ADDR adjusted_pc
+ = pc - gdbarch_decr_pc_after_break (gdbarch);
+
+ if (adjusted_pc != pc)
+ regcache_write_pc (regcache, adjusted_pc);
+ }
+ }
+ else
+ error (_("wave_get_info for wave_%ld failed (%s)"),
+ wave_id.handle, get_status_string (status));
+
+ info->wave_events.emplace_back (event_ptid, ws);
+ break;
+ }
+
+ case AMD_DBGAPI_EVENT_KIND_CODE_OBJECT_LIST_UPDATED:
+ /* We get here when the following sequence of events happens:
+
+ - the inferior hits the amd-dbgapi "r_brk" internal breakpoint
+ - amd_dbgapi_target_breakpoint::check_status calls
+ amd_dbgapi_report_breakpoint_hit, which queues an event of this
+ kind in dbgapi
+ - amd_dbgapi_target_breakpoint::check_status calls
+ process_event_queue, which pulls the event out of dbgapi, and
+ gets us here
+
+ When amd_dbgapi_target_breakpoint::check_status is called, the current
+ inferior is the inferior that hit the breakpoint, which should still be
+ the case now. */
+ gdb_assert (inf == current_inferior ());
+ handle_solib_event ();
+ break;
+
+ case AMD_DBGAPI_EVENT_KIND_BREAKPOINT_RESUME:
+ /* Breakpoint resume events should be handled by the breakpoint
+ action, and this code should not reach this. */
+ gdb_assert_not_reached ("unhandled event kind");
+ break;
+
+ case AMD_DBGAPI_EVENT_KIND_RUNTIME:
+ {
+ amd_dbgapi_runtime_state_t runtime_state;
+
+ status = amd_dbgapi_event_get_info (event_id,
+ AMD_DBGAPI_EVENT_INFO_RUNTIME_STATE,
+ sizeof (runtime_state),
+ &runtime_state);
+ if (status != AMD_DBGAPI_STATUS_SUCCESS)
+ error (_("event_get_info for event_%ld failed (%s)"),
+ event_id.handle, get_status_string (status));
+
+ gdb_assert (runtime_state == AMD_DBGAPI_RUNTIME_STATE_UNLOADED);
+ gdb_assert
+ (info->runtime_state == AMD_DBGAPI_RUNTIME_STATE_LOADED_SUCCESS);
+
+ info->runtime_state = runtime_state;
+
+ gdb_assert (inf->target_is_pushed (&the_amd_dbgapi_target));
+ inf->unpush_target (&the_amd_dbgapi_target);
+ }
+ break;
+
+ default:
+ error (_("event kind (%d) not supported"), event_kind);
+ }
+}
+
+/* Return a textual version of KIND. */
+
+static const char *
+event_kind_str (amd_dbgapi_event_kind_t kind)
+{
+ switch (kind)
+ {
+ case AMD_DBGAPI_EVENT_KIND_NONE:
+ return "NONE";
+
+ case AMD_DBGAPI_EVENT_KIND_WAVE_STOP:
+ return "WAVE_STOP";
+
+ case AMD_DBGAPI_EVENT_KIND_WAVE_COMMAND_TERMINATED:
+ return "WAVE_COMMAND_TERMINATED";
+
+ case AMD_DBGAPI_EVENT_KIND_CODE_OBJECT_LIST_UPDATED:
+ return "CODE_OBJECT_LIST_UPDATED";
+
+ case AMD_DBGAPI_EVENT_KIND_BREAKPOINT_RESUME:
+ return "BREAKPOINT_RESUME";
+
+ case AMD_DBGAPI_EVENT_KIND_RUNTIME:
+ return "RUNTIME";
+
+ case AMD_DBGAPI_EVENT_KIND_QUEUE_ERROR:
+ return "QUEUE_ERROR";
+ }
+
+ gdb_assert_not_reached ("unhandled amd_dbgapi_event_kind_t value");
+}
+
+/* Drain the dbgapi event queue of a given process_id, or of all processes if
+ process_id is AMD_DBGAPI_PROCESS_NONE. Stop processing the events if an
+ event of a given kind is requested and `process_id` is not
+ AMD_DBGAPI_PROCESS_NONE. Wave stop events that are not returned are queued
+ into their inferior's amd_dbgapi_inferior_info pending wave events. */
+
+static amd_dbgapi_event_id_t
+process_event_queue (amd_dbgapi_process_id_t process_id,
+ amd_dbgapi_event_kind_t until_event_kind)
+{
+ /* An event of a given type can only be requested from a single
+ process_id. */
+ gdb_assert (until_event_kind == AMD_DBGAPI_EVENT_KIND_NONE
+ || process_id != AMD_DBGAPI_PROCESS_NONE);
+
+ while (true)
+ {
+ amd_dbgapi_event_id_t event_id;
+ amd_dbgapi_event_kind_t event_kind;
+
+ amd_dbgapi_status_t status
+ = amd_dbgapi_process_next_pending_event (process_id, &event_id,
+ &event_kind);
+
+ if (status != AMD_DBGAPI_STATUS_SUCCESS)
+ error (_("next_pending_event failed (%s)"), get_status_string (status));
+
+ if (event_kind != AMD_DBGAPI_EVENT_KIND_NONE)
+ amd_dbgapi_debug_printf ("Pulled event from dbgapi: "
+ "event_id.handle = %" PRIu64 ", "
+ "event_kind = %s",
+ event_id.handle,
+ event_kind_str (event_kind));
+
+ if (event_id == AMD_DBGAPI_EVENT_NONE || event_kind == until_event_kind)
+ return event_id;
+
+ process_one_event (event_id, event_kind);
+ }
+}
+
+bool
+amd_dbgapi_target::has_pending_events ()
+{
+ if (amd_dbgapi_async_event_handler != nullptr
+ && async_event_handler_marked (amd_dbgapi_async_event_handler))
+ return true;
+
+ return beneath ()->has_pending_events ();
+}
+
+/* Pop one pending event from the per-inferior structures.
+
+ If PID is not -1, restrict the search to the inferior with that pid. */
+
+static std::pair<ptid_t, target_waitstatus>
+consume_one_event (int pid)
+{
+ auto *target = current_inferior ()->process_target ();
+ struct amd_dbgapi_inferior_info *info = nullptr;
+
+ if (pid == -1)
+ {
+ for (inferior *inf : all_inferiors (target))
+ {
+ info = get_amd_dbgapi_inferior_info (inf);
+ if (!info->wave_events.empty ())
+ break;
+ }
+
+ gdb_assert (info != nullptr);
+ }
+ else
+ {
+ inferior *inf = find_inferior_pid (target, pid);
+
+ gdb_assert (inf != nullptr);
+ info = get_amd_dbgapi_inferior_info (inf);
+ }
+
+ if (info->wave_events.empty ())
+ return { minus_one_ptid, {} };
+
+ auto event = info->wave_events.front ();
+ info->wave_events.pop_front ();
+
+ return event;
+}
+
+ptid_t
+amd_dbgapi_target::wait (ptid_t ptid, struct target_waitstatus *ws,
+ target_wait_flags target_options)
+{
+ gdb_assert (!current_inferior ()->process_target ()->commit_resumed_state);
+ gdb_assert (ptid == minus_one_ptid || ptid.is_pid ());
+
+ amd_dbgapi_debug_printf ("ptid = %s", ptid.to_string ().c_str ());
+
+ ptid_t event_ptid = beneath ()->wait (ptid, ws, target_options);
+ if (event_ptid != minus_one_ptid)
+ {
+ if (ws->kind () == TARGET_WAITKIND_EXITED
+ || ws->kind () == TARGET_WAITKIND_SIGNALLED)
+ {
+ /* This inferior has exited so drain its dbgapi event queue. */
+ while (consume_one_event (event_ptid.pid ()).first
+ != minus_one_ptid)
+ ;
+ }
+ return event_ptid;
+ }
+
+ gdb_assert (ws->kind () == TARGET_WAITKIND_NO_RESUMED
+ || ws->kind () == TARGET_WAITKIND_IGNORE);
+
+ /* Flush the async handler first. */
+ if (target_is_async_p ())
+ async_event_handler_clear ();
+
+ /* There may be more events to process (either already in `wave_events` or
+ that we need to fetch from dbgapi. Mark the async event handler so that
+ amd_dbgapi_target::wait gets called again and again, until it eventually
+ returns minus_one_ptid. */
+ auto more_events = make_scope_exit ([] ()
+ {
+ if (target_is_async_p ())
+ async_event_handler_mark ();
+ });
+
+ auto *proc_target = current_inferior ()->process_target ();
+
+ /* Disable forward progress for the specified pid in ptid if it isn't
+ minus_on_ptid, or all attached processes if ptid is minus_one_ptid. */
+ require_forward_progress (ptid, proc_target, false);
+
+ target_waitstatus gpu_waitstatus;
+ std::tie (event_ptid, gpu_waitstatus) = consume_one_event (ptid.pid ());
+ if (event_ptid == minus_one_ptid)
+ {
+ /* Drain the events from the amd_dbgapi and preserve the ordering. */
+ process_event_queue ();
+
+ std::tie (event_ptid, gpu_waitstatus) = consume_one_event (ptid.pid ());
+ if (event_ptid == minus_one_ptid)
+ {
+ /* If we requested a specific ptid, and nothing came out, assume
+ another ptid may have more events, otherwise, keep the
+ async_event_handler flushed. */
+ if (ptid == minus_one_ptid)
+ more_events.release ();
+
+ if (ws->kind () == TARGET_WAITKIND_NO_RESUMED)
+ {
+ /* We can't easily check that all GPU waves are stopped, and no
+ new waves can be created (the GPU has fixed function hardware
+ to create new threads), so even if the target beneath returns
+ waitkind_no_resumed, we have to report waitkind_ignore if GPU
+ debugging is enabled for at least one resumed inferior handled
+ by the amd-dbgapi target. */
+
+ for (inferior *inf : all_inferiors ())
+ if (inf->target_at (arch_stratum) == &the_amd_dbgapi_target
+ && get_amd_dbgapi_inferior_info (inf)->runtime_state
+ == AMD_DBGAPI_RUNTIME_STATE_LOADED_SUCCESS)
+ {
+ ws->set_ignore ();
+ break;
+ }
+ }
+
+ /* There are no events to report, return the target beneath's
+ waitstatus (either IGNORE or NO_RESUMED). */
+ return minus_one_ptid;
+ }
+ }
+
+ *ws = gpu_waitstatus;
+ return event_ptid;
+}
+
+bool
+amd_dbgapi_target::stopped_by_sw_breakpoint ()
+{
+ if (!ptid_is_gpu (inferior_ptid))
+ return beneath ()->stopped_by_sw_breakpoint ();
+
+ amd_dbgapi_wave_id_t wave_id = get_amd_dbgapi_wave_id (inferior_ptid);
+
+ amd_dbgapi_wave_stop_reasons_t stop_reason;
+ amd_dbgapi_status_t status
+ = amd_dbgapi_wave_get_info (wave_id, AMD_DBGAPI_WAVE_INFO_STOP_REASON,
+ sizeof (stop_reason), &stop_reason);
+
+ if (status != AMD_DBGAPI_STATUS_SUCCESS)
+ return false;
+
+ return (stop_reason & AMD_DBGAPI_WAVE_STOP_REASON_BREAKPOINT) != 0;
+}
+
+bool
+amd_dbgapi_target::stopped_by_hw_breakpoint ()
+{
+ if (!ptid_is_gpu (inferior_ptid))
+ return beneath ()->stopped_by_hw_breakpoint ();
+
+ return false;
+}
+
+/* Make the amd-dbgapi library attach to the process behind INF.
+
+ Note that this is unrelated to the "attach" GDB concept / command.
+
+ By attaching to the process, we get a notifier fd that tells us when it
+ activates the ROCm runtime and when there are subsequent debug events. */
+
+static void
+attach_amd_dbgapi (inferior *inf)
+{
+ AMD_DBGAPI_SCOPED_DEBUG_START_END ("inf num = %d", inf->num);
+
+ if (!target_can_async_p ())
+ {
+ warning (_("The amd-dbgapi target requires the target beneath to be "
+ "asynchronous, GPU debugging is disabled"));
+ return;
+ }
+
+ auto *info = get_amd_dbgapi_inferior_info (inf);
+
+ /* Are we already attached? */
+ if (info->process_id != AMD_DBGAPI_PROCESS_NONE)
+ {
+ amd_dbgapi_debug_printf
+ ("already attached: process_id = %" PRIu64, info->process_id.handle);
+ return;
+ }
+
+ amd_dbgapi_status_t status
+ = amd_dbgapi_process_attach
+ (reinterpret_cast<amd_dbgapi_client_process_id_t> (inf),
+ &info->process_id);
+ if (status == AMD_DBGAPI_STATUS_ERROR_RESTRICTION)
+ {
+ warning (_("amd-dbgapi: unable to enable GPU debugging due to a "
+ "restriction error"));
+ return;
+ }
+ else if (status != AMD_DBGAPI_STATUS_SUCCESS)
+ {
+ warning (_("amd-dbgapi: could not attach to process %d (%s), GPU "
+ "debugging will not be available."), inf->pid,
+ get_status_string (status));
+ return;
+ }
+
+ if (amd_dbgapi_process_get_info (info->process_id,
+ AMD_DBGAPI_PROCESS_INFO_NOTIFIER,
+ sizeof (info->notifier), &info->notifier)
+ != AMD_DBGAPI_STATUS_SUCCESS)
+ {
+ amd_dbgapi_process_detach (info->process_id);
+ info->process_id = AMD_DBGAPI_PROCESS_NONE;
+ warning (_("amd-dbgapi: could not retrieve process %d's notifier, GPU "
+ "debugging will not be available."), inf->pid);
+ return;
+ }
+
+ amd_dbgapi_debug_printf ("process_id = %" PRIu64 ", notifier fd = %d",
+ info->process_id.handle, info->notifier);
+
+ /* 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. */
+ dbgapi_notifier_handler (0, info);
+
+ add_file_handler (info->notifier, dbgapi_notifier_handler, info,
+ "amd-dbgapi notifier");
+}
+
+static void maybe_reset_amd_dbgapi ();
+
+/* Make the amd-dbgapi library detach from INF.
+
+ Note that this us unrelated to the "detach" GDB concept / command.
+
+ This undoes what attach_amd_dbgapi does. */
+
+static void
+detach_amd_dbgapi (inferior *inf)
+{
+ AMD_DBGAPI_SCOPED_DEBUG_START_END ("inf num = %d", inf->num);
+
+ auto *info = get_amd_dbgapi_inferior_info (inf);
+
+ if (info->process_id == AMD_DBGAPI_PROCESS_NONE)
+ return;
+
+ info->runtime_state = AMD_DBGAPI_RUNTIME_STATE_UNLOADED;
+
+ amd_dbgapi_status_t status = amd_dbgapi_process_detach (info->process_id);
+ if (status != AMD_DBGAPI_STATUS_SUCCESS)
+ warning (_("amd-dbgapi: could not detach from process %d (%s)"),
+ inf->pid, get_status_string (status));
+
+ gdb_assert (info->notifier != -1);
+ delete_file_handler (info->notifier);
+
+ /* This is a noop if the target is not pushed. */
+ inf->unpush_target (&the_amd_dbgapi_target);
+
+ /* Delete the breakpoints that are still active. */
+ for (auto &&value : info->breakpoint_map)
+ delete_breakpoint (value.second);
+
+ /* Reset the amd_dbgapi_inferior_info. */
+ *info = amd_dbgapi_inferior_info (inf);
+
+ maybe_reset_amd_dbgapi ();
+}
+
+void
+amd_dbgapi_target::mourn_inferior ()
+{
+ detach_amd_dbgapi (current_inferior ());
+ beneath ()->mourn_inferior ();
+}
+
+void
+amd_dbgapi_target::detach (inferior *inf, int from_tty)
+{
+ /* We're about to resume the waves by detaching the dbgapi library from the
+ inferior, so we need to remove all breakpoints that are still inserted.
+
+ Breakpoints may still be inserted because the inferior may be running in
+ non-stop mode, or because GDB changed the default setting to leave all
+ breakpoints inserted in all-stop mode when all threads are stopped. */
+ remove_breakpoints_inf (current_inferior ());
+
+ detach_amd_dbgapi (inf);
+ beneath ()->detach (inf, from_tty);
+}
+
+void
+amd_dbgapi_target::fetch_registers (struct regcache *regcache, int regno)
+{
+ if (!ptid_is_gpu (regcache->ptid ()))
+ {
+ beneath ()->fetch_registers (regcache, regno);
+ return;
+ }
+
+ struct gdbarch *gdbarch = regcache->arch ();
+ gdb_assert (is_amdgpu_arch (gdbarch));
+
+ amdgpu_gdbarch_tdep *tdep = get_amdgpu_gdbarch_tdep (gdbarch);
+ amd_dbgapi_wave_id_t wave_id = get_amd_dbgapi_wave_id (regcache->ptid ());
+ gdb_byte raw[AMDGPU_MAX_REGISTER_SIZE];
+ amd_dbgapi_status_t status
+ = amd_dbgapi_read_register (wave_id, tdep->register_ids[regno], 0,
+ register_type (gdbarch, regno)->length (),
+ raw);
+
+ if (status == AMD_DBGAPI_STATUS_SUCCESS)
+ regcache->raw_supply (regno, raw);
+ else if (status != AMD_DBGAPI_STATUS_ERROR_REGISTER_NOT_AVAILABLE)
+ warning (_("Couldn't read register %s (#%d) (%s)."),
+ gdbarch_register_name (gdbarch, regno), regno,
+ get_status_string (status));
+}
+
+void
+amd_dbgapi_target::store_registers (struct regcache *regcache, int regno)
+{
+ if (!ptid_is_gpu (regcache->ptid ()))
+ {
+ beneath ()->store_registers (regcache, regno);
+ return;
+ }
+
+ struct gdbarch *gdbarch = regcache->arch ();
+ gdb_assert (is_amdgpu_arch (gdbarch));
+
+ gdb_byte raw[AMDGPU_MAX_REGISTER_SIZE];
+ regcache->raw_collect (regno, &raw);
+
+ amdgpu_gdbarch_tdep *tdep = get_amdgpu_gdbarch_tdep (gdbarch);
+
+ /* If the register has read-only bits, invalidate the value in the regcache
+ as the value actualy written may differ. */
+ if (tdep->register_properties[regno]
+ & AMD_DBGAPI_REGISTER_PROPERTY_READONLY_BITS)
+ regcache->invalidate (regno);
+
+ /* Invalidate all volatile registers if this register has the invalidate
+ volatile property. For example, writting to VCC may change the content
+ of STATUS.VCCZ. */
+ if (tdep->register_properties[regno]
+ & AMD_DBGAPI_REGISTER_PROPERTY_INVALIDATE_VOLATILE)
+ {
+ for (size_t r = 0; r < tdep->register_properties.size (); ++r)
+ if (tdep->register_properties[r] & AMD_DBGAPI_REGISTER_PROPERTY_VOLATILE)
+ regcache->invalidate (r);
+ }
+
+ amd_dbgapi_wave_id_t wave_id = get_amd_dbgapi_wave_id (regcache->ptid ());
+ amd_dbgapi_status_t status
+ = amd_dbgapi_write_register (wave_id, tdep->register_ids[regno], 0,
+ register_type (gdbarch, regno)->length (),
+ raw);
+
+ if (status != AMD_DBGAPI_STATUS_SUCCESS)
+ warning (_("Couldn't write register %s (#%d)."),
+ gdbarch_register_name (gdbarch, regno), regno);
+}
+
+struct gdbarch *
+amd_dbgapi_target::thread_architecture (ptid_t ptid)
+{
+ if (!ptid_is_gpu (ptid))
+ return beneath ()->thread_architecture (ptid);
+
+ /* We can cache the gdbarch for a given wave_id (ptid::tid) because
+ wave IDs are unique, and aren't reused. */
+ if (ptid.tid () == m_cached_arch_tid)
+ return m_cached_arch;
+
+ amd_dbgapi_wave_id_t wave_id = get_amd_dbgapi_wave_id (ptid);
+ amd_dbgapi_architecture_id_t architecture_id;
+ amd_dbgapi_status_t status;
+
+ status = amd_dbgapi_wave_get_info (wave_id, AMD_DBGAPI_WAVE_INFO_ARCHITECTURE,
+ sizeof (architecture_id),
+ &architecture_id);
+ if (status != AMD_DBGAPI_STATUS_SUCCESS)
+ error (_("Couldn't get architecture for wave_%ld"), ptid.tid ());
+
+ uint32_t elf_amdgpu_machine;
+ status = amd_dbgapi_architecture_get_info
+ (architecture_id, AMD_DBGAPI_ARCHITECTURE_INFO_ELF_AMDGPU_MACHINE,
+ sizeof (elf_amdgpu_machine), &elf_amdgpu_machine);
+ if (status != AMD_DBGAPI_STATUS_SUCCESS)
+ error (_("Couldn't get elf_amdgpu_machine for architecture_%ld"),
+ architecture_id.handle);
+
+ struct gdbarch_info info;
+ info.bfd_arch_info = bfd_lookup_arch (bfd_arch_amdgcn, elf_amdgpu_machine);
+ info.byte_order = BFD_ENDIAN_LITTLE;
+
+ m_cached_arch_tid = ptid.tid ();
+ m_cached_arch = gdbarch_find_by_info (info);
+ if (m_cached_arch == nullptr)
+ error (_("Couldn't get elf_amdgpu_machine (%#x)"), elf_amdgpu_machine);
+
+ return m_cached_arch;
+}
+
+void
+amd_dbgapi_target::thread_events (int enable)
+{
+ m_report_thread_events = enable;
+ beneath ()->thread_events (enable);
+}
+
+void
+amd_dbgapi_target::update_thread_list ()
+{
+ for (inferior *inf : all_inferiors ())
+ {
+ amd_dbgapi_process_id_t process_id
+ = get_amd_dbgapi_process_id (inf);
+ if (process_id == AMD_DBGAPI_PROCESS_NONE)
+ {
+ /* The inferior may not be attached yet. */
+ continue;
+ }
+
+ size_t count;
+ amd_dbgapi_wave_id_t *wave_list;
+ amd_dbgapi_changed_t changed;
+ amd_dbgapi_status_t status
+ = amd_dbgapi_process_wave_list (process_id, &count, &wave_list,
+ &changed);
+ if (status != AMD_DBGAPI_STATUS_SUCCESS)
+ error (_("amd_dbgapi_wave_list failed (%s)"),
+ get_status_string (status));
+
+ if (changed == AMD_DBGAPI_CHANGED_NO)
+ continue;
+
+ /* Create a set and free the wave list. */
+ std::set<ptid_t::tid_type> threads;
+ for (size_t i = 0; i < count; ++i)
+ threads.emplace (wave_list[i].handle);
+
+ xfree (wave_list);
+
+ /* Prune the wave_ids that already have a thread_info. Any thread_info
+ which does not have a corresponding wave_id represents a wave which
+ is gone at this point and should be deleted. */
+ for (thread_info *tp : inf->threads_safe ())
+ if (ptid_is_gpu (tp->ptid) && tp->state != THREAD_EXITED)
+ {
+ auto it = threads.find (tp->ptid.tid ());
+
+ if (it == threads.end ())
+ delete_thread (tp);
+ else
+ threads.erase (it);
+ }
+
+ /* The wave_ids that are left require a new thread_info. */
+ for (ptid_t::tid_type tid : threads)
+ {
+ ptid_t wave_ptid
+ = make_gpu_ptid (inf->pid, amd_dbgapi_wave_id_t {tid});
+
+ add_thread_silent (inf->process_target (), wave_ptid);
+ set_running (inf->process_target (), wave_ptid, true);
+ set_executing (inf->process_target (), wave_ptid, true);
+ }
+ }
+
+ /* Give the beneath target a chance to do extra processing. */
+ this->beneath ()->update_thread_list ();
+}
+
+/* inferior_created observer. */
+
+static void
+amd_dbgapi_target_inferior_created (inferior *inf)
+{
+ /* If the inferior is not running on the native target (e.g. it is running
+ on a remote target), we don't want to deal with it. */
+ if (inf->process_target () != get_native_target ())
+ return;
+
+ attach_amd_dbgapi (inf);
+}
+
+/* inferior_exit observer.
+
+ This covers normal exits, but also detached inferiors (including detached
+ fork parents). */
+
+static void
+amd_dbgapi_inferior_exited (inferior *inf)
+{
+ detach_amd_dbgapi (inf);
+}
+
+/* inferior_pre_detach observer. */
+
+static void
+amd_dbgapi_inferior_pre_detach (inferior *inf)
+{
+ /* We need to amd-dbgapi-detach before we ptrace-detach. If the amd-dbgapi
+ target isn't pushed, do that now. If the amd-dbgapi target is pushed,
+ we'll do it in amd_dbgapi_target::detach. */
+ if (!inf->target_is_pushed (&the_amd_dbgapi_target))
+ detach_amd_dbgapi (inf);
+}
+
+/* get_os_pid callback. */
+
+static amd_dbgapi_status_t
+amd_dbgapi_get_os_pid_callback
+ (amd_dbgapi_client_process_id_t client_process_id, pid_t *pid)
+{
+ inferior *inf = reinterpret_cast<inferior *> (client_process_id);
+
+ if (inf->pid == 0)
+ return AMD_DBGAPI_STATUS_ERROR_PROCESS_EXITED;
+
+ *pid = inf->pid;
+ return AMD_DBGAPI_STATUS_SUCCESS;
+}
+
+/* insert_breakpoint callback. */
+
+static amd_dbgapi_status_t
+amd_dbgapi_insert_breakpoint_callback
+ (amd_dbgapi_client_process_id_t client_process_id,
+ amd_dbgapi_global_address_t address,
+ amd_dbgapi_breakpoint_id_t breakpoint_id)
+{
+ inferior *inf = reinterpret_cast<inferior *> (client_process_id);
+ struct amd_dbgapi_inferior_info *info = get_amd_dbgapi_inferior_info (inf);
+
+ auto it = info->breakpoint_map.find (breakpoint_id.handle);
+ if (it != info->breakpoint_map.end ())
+ return AMD_DBGAPI_STATUS_ERROR_INVALID_BREAKPOINT_ID;
+
+ /* We need to find the address in the given inferior's program space. */
+ scoped_restore_current_thread restore_thread;
+ switch_to_inferior_no_thread (inf);
+
+ /* Create a new breakpoint. */
+ struct obj_section *section = find_pc_section (address);
+ if (section == nullptr || section->objfile == nullptr)
+ return AMD_DBGAPI_STATUS_ERROR;
+
+ std::unique_ptr<breakpoint> bp_up
+ (new amd_dbgapi_target_breakpoint (section->objfile->arch (), address));
+
+ breakpoint *bp = install_breakpoint (true, std::move (bp_up), 1);
+
+ info->breakpoint_map.emplace (breakpoint_id.handle, bp);
+ return AMD_DBGAPI_STATUS_SUCCESS;
+}
+
+/* remove_breakpoint callback. */
+
+static amd_dbgapi_status_t
+amd_dbgapi_remove_breakpoint_callback
+ (amd_dbgapi_client_process_id_t client_process_id,
+ amd_dbgapi_breakpoint_id_t breakpoint_id)
+{
+ inferior *inf = reinterpret_cast<inferior *> (client_process_id);
+ struct amd_dbgapi_inferior_info *info = get_amd_dbgapi_inferior_info (inf);
+
+ auto it = info->breakpoint_map.find (breakpoint_id.handle);
+ if (it == info->breakpoint_map.end ())
+ return AMD_DBGAPI_STATUS_ERROR_INVALID_BREAKPOINT_ID;
+
+ delete_breakpoint (it->second);
+ info->breakpoint_map.erase (it);
+
+ return AMD_DBGAPI_STATUS_SUCCESS;
+}
+
+/* Style for some kinds of messages. */
+
+static cli_style_option fatal_error_style
+ ("amd_dbgapi_fatal_error", ui_file_style::RED);
+static cli_style_option warning_style
+ ("amd_dbgapi_warning", ui_file_style::YELLOW);
+
+/* BLACK + BOLD means dark gray. */
+static cli_style_option trace_style
+ ("amd_dbgapi_trace", ui_file_style::BLACK, ui_file_style::BOLD);
+
+/* log_message callback. */
+
+static void
+amd_dbgapi_log_message_callback (amd_dbgapi_log_level_t level,
+ const char *message)
+{
+ gdb::optional<target_terminal::scoped_restore_terminal_state> tstate;
+
+ if (target_supports_terminal_ours ())
+ {
+ tstate.emplace ();
+ target_terminal::ours_for_output ();
+ }
+
+ /* Error and warning messages are meant to be printed to the user. */
+ if (level == AMD_DBGAPI_LOG_LEVEL_FATAL_ERROR
+ || level == AMD_DBGAPI_LOG_LEVEL_WARNING)
+ {
+ begin_line ();
+ ui_file_style style = (level == AMD_DBGAPI_LOG_LEVEL_FATAL_ERROR
+ ? fatal_error_style : warning_style).style ();
+ gdb_printf (gdb_stderr, "%ps\n", styled_string (style, message));
+ return;
+ }
+
+ /* Print other messages as debug logs. TRACE and VERBOSE messages are
+ very verbose, print them dark grey so it's easier to spot other messages
+ through the flood. */
+ if (level >= AMD_DBGAPI_LOG_LEVEL_TRACE)
+ {
+ debug_prefixed_printf (amd_dbgapi_lib_debug_module (), nullptr, "%ps",
+ styled_string (trace_style.style (), message));
+ return;
+ }
+
+ debug_prefixed_printf (amd_dbgapi_lib_debug_module (), nullptr, "%s",
+ message);
+}
+
+/* Callbacks passed to amd_dbgapi_initialize. */
+
+static amd_dbgapi_callbacks_t dbgapi_callbacks = {
+ .allocate_memory = malloc,
+ .deallocate_memory = free,
+ .get_os_pid = amd_dbgapi_get_os_pid_callback,
+ .insert_breakpoint = amd_dbgapi_insert_breakpoint_callback,
+ .remove_breakpoint = amd_dbgapi_remove_breakpoint_callback,
+ .log_message = amd_dbgapi_log_message_callback,
+};
+
+void
+amd_dbgapi_target::close ()
+{
+ if (amd_dbgapi_async_event_handler != nullptr)
+ delete_async_event_handler (&amd_dbgapi_async_event_handler);
+}
+
+/* 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;
+
+/* Mapping from amd-dbgapi log level enum values to text. */
+
+static constexpr const char *debug_amd_dbgapi_lib_log_level_enums[] =
+{
+ /* [AMD_DBGAPI_LOG_LEVEL_NONE] = */ "off",
+ /* [AMD_DBGAPI_LOG_LEVEL_FATAL_ERROR] = */ "error",
+ /* [AMD_DBGAPI_LOG_LEVEL_WARNING] = */ "warning",
+ /* [AMD_DBGAPI_LOG_LEVEL_INFO] = */ "info",
+ /* [AMD_DBGAPI_LOG_LEVEL_TRACE] = */ "trace",
+ /* [AMD_DBGAPI_LOG_LEVEL_VERBOSE] = */ "verbose",
+ nullptr
+};
+
+/* Storage for "set debug amd-dbgapi-lib log-level". */
+
+static const char *debug_amd_dbgapi_lib_log_level
+ = debug_amd_dbgapi_lib_log_level_enums[AMD_DBGAPI_LOG_LEVEL_WARNING];
+
+/* Get the amd-dbgapi library log level requested by the user. */
+
+static amd_dbgapi_log_level_t
+get_debug_amd_dbgapi_lib_log_level ()
+{
+ for (size_t pos = 0;
+ debug_amd_dbgapi_lib_log_level_enums[pos] != nullptr;
+ ++pos)
+ if (debug_amd_dbgapi_lib_log_level
+ == debug_amd_dbgapi_lib_log_level_enums[pos])
+ return static_cast<amd_dbgapi_log_level_t> (pos);
+
+ gdb_assert_not_reached ("invalid log level");
+}
+
+/* Callback for "set debug amd-dbgapi log-level", apply the selected log level
+ to the library. */
+
+static void
+set_debug_amd_dbgapi_lib_log_level (const char *args, int from_tty,
+ struct cmd_list_element *c)
+{
+ amd_dbgapi_set_log_level (get_debug_amd_dbgapi_lib_log_level ());
+}
+
+/* Callback for "show debug amd-dbgapi log-level". */
+
+static void
+show_debug_amd_dbgapi_lib_log_level (struct ui_file *file, int from_tty,
+ struct cmd_list_element *c,
+ const char *value)
+{
+ gdb_printf (file, _("The amd-dbgapi library log level is %s.\n"), value);
+}
+
+/* If the amd-dbgapi library is not attached to any process, finalize and
+ re-initialize it so that the handle ID numbers will all start from the
+ beginning again. This is only for convenience, not essential. */
+
+static void
+maybe_reset_amd_dbgapi ()
+{
+ for (inferior *inf : all_non_exited_inferiors ())
+ {
+ amd_dbgapi_inferior_info *info = get_amd_dbgapi_inferior_info (inf);
+
+ if (info->process_id != AMD_DBGAPI_PROCESS_NONE)
+ return;
+ }
+
+ amd_dbgapi_status_t status = amd_dbgapi_finalize ();
+ if (status != AMD_DBGAPI_STATUS_SUCCESS)
+ error (_("amd-dbgapi failed to finalize (%s)"),
+ get_status_string (status));
+
+ status = amd_dbgapi_initialize (&dbgapi_callbacks);
+ if (status != AMD_DBGAPI_STATUS_SUCCESS)
+ error (_("amd-dbgapi failed to initialize (%s)"),
+ get_status_string (status));
+}
+
+extern initialize_file_ftype _initialize_amd_dbgapi_target;
+
+void
+_initialize_amd_dbgapi_target ()
+{
+ /* Make sure the loaded debugger library version is greater than or equal to
+ the one used to build GDB. */
+ uint32_t major, minor, patch;
+ amd_dbgapi_get_version (&major, &minor, &patch);
+ if (major != AMD_DBGAPI_VERSION_MAJOR || minor < AMD_DBGAPI_VERSION_MINOR)
+ error (_("amd-dbgapi library version mismatch, got %d.%d.%d, need %d.%d+"),
+ major, minor, patch, AMD_DBGAPI_VERSION_MAJOR,
+ AMD_DBGAPI_VERSION_MINOR);
+
+ /* Initialize the AMD Debugger API. */
+ amd_dbgapi_status_t status = amd_dbgapi_initialize (&dbgapi_callbacks);
+ if (status != AMD_DBGAPI_STATUS_SUCCESS)
+ error (_("amd-dbgapi failed to initialize (%s)"),
+ get_status_string (status));
+
+ /* Set the initial log level. */
+ amd_dbgapi_set_log_level (get_debug_amd_dbgapi_lib_log_level ());
+
+ /* Install observers. */
+ gdb::observers::inferior_created.attach
+ (amd_dbgapi_target_inferior_created,
+ amd_dbgapi_target_inferior_created_observer_token, "amd-dbgapi");
+ 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 ("amd-dbgapi-lib", no_class,
+ _("Generic command for setting amd-dbgapi library "
+ "debugging flags."),
+ &set_debug_amd_dbgapi_lib_list, 0, &setdebuglist);
+
+ add_show_prefix_cmd ("amd-dbgapi-lib", no_class,
+ _("Generic command for showing amd-dbgapi library "
+ "debugging flags."),
+ &show_debug_amd_dbgapi_lib_list, 0, &showdebuglist);
+
+ add_setshow_enum_cmd ("log-level", class_maintenance,
+ debug_amd_dbgapi_lib_log_level_enums,
+ &debug_amd_dbgapi_lib_log_level,
+ _("Set the amd-dbgapi library log level."),
+ _("Show the amd-dbgapi library log level."),
+ _("off == no logging is enabled\n"
+ "error == fatal errors are reported\n"
+ "warning == fatal errors and warnings are reported\n"
+ "info == fatal errors, warnings, and info "
+ "messages are reported\n"
+ "trace == fatal errors, warnings, info, and "
+ "API tracing messages are reported\n"
+ "verbose == all messages are reported"),
+ set_debug_amd_dbgapi_lib_log_level,
+ show_debug_amd_dbgapi_lib_log_level,
+ &set_debug_amd_dbgapi_lib_list,
+ &show_debug_amd_dbgapi_lib_list);
+
+ add_setshow_boolean_cmd ("amd-dbgapi", class_maintenance,
+ &debug_amd_dbgapi,
+ _("Set debugging of amd-dbgapi target."),
+ _("Show debugging of amd-dbgapi target."),
+ _("\
+When on, print debug messages relating to the amd-dbgapi target."),
+ nullptr, nullptr,
+ &setdebuglist, &showdebuglist);
+}
diff --git a/gdb/amd-dbgapi-target.h b/gdb/amd-dbgapi-target.h
new file mode 100644
index 00000000000..beff2ad0bed
--- /dev/null
+++ b/gdb/amd-dbgapi-target.h
@@ -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 */
diff --git a/gdb/amdgpu-tdep.c b/gdb/amdgpu-tdep.c
new file mode 100644
index 00000000000..fc5e2438c7f
--- /dev/null
+++ b/gdb/amdgpu-tdep.c
@@ -0,0 +1,1367 @@
+/* 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/>. */
+
+#include "defs.h"
+
+#include "amd-dbgapi-target.h"
+#include "amdgpu-tdep.h"
+#include "arch-utils.h"
+#include "disasm.h"
+#include "dwarf2/frame.h"
+#include "frame-unwind.h"
+#include "gdbarch.h"
+#include "gdbsupport/selftest.h"
+#include "gdbtypes.h"
+#include "inferior.h"
+#include "objfiles.h"
+#include "observable.h"
+#include "producer.h"
+#include "reggroups.h"
+
+/* See amdgpu-tdep.h. */
+
+bool
+is_amdgpu_arch (struct gdbarch *arch)
+{
+ gdb_assert (arch != nullptr);
+ return gdbarch_bfd_arch_info (arch)->arch == bfd_arch_amdgcn;
+}
+
+/* See amdgpu-tdep.h. */
+
+amdgpu_gdbarch_tdep *
+get_amdgpu_gdbarch_tdep (gdbarch *arch)
+{
+ return gdbarch_tdep<amdgpu_gdbarch_tdep> (arch);
+}
+
+/* Return the name of register REGNUM. */
+
+static const char *
+amdgpu_register_name (struct gdbarch *gdbarch, int regnum)
+{
+ /* The list of registers reported by amd-dbgapi for a given architecture
+ contains some duplicate names. For instance, there is an "exec" register
+ for waves in the wave32 mode and one for the waves in the wave64 mode.
+ However, at most one register with a given name is actually allocated for
+ a specific wave. If INFERIOR_PTID represents a GPU wave, we query
+ amd-dbgapi to know whether the requested register actually exists for the
+ current wave, so there won't be duplicates in the the register names we
+ report for that wave.
+
+ But there are two known cases where INFERIOR_PTID doesn't represent a GPU
+ wave:
+
+ - The user does "set arch amdgcn:gfxNNN" followed with "maint print
+ registers"
+ - The "register_name" selftest
+
+ In these cases, we can't query amd-dbgapi to know whether we should hide
+ the register or not. The "register_name" selftest checks that there aren't
+ duplicates in the register names returned by the gdbarch, so if we simply
+ return all register names, that test will fail. The other simple option is
+ to never return a register name, which is what we do here. */
+ if (!ptid_is_gpu (inferior_ptid))
+ return "";
+
+ amd_dbgapi_wave_id_t wave_id = get_amd_dbgapi_wave_id (inferior_ptid);
+ amdgpu_gdbarch_tdep *tdep = get_amdgpu_gdbarch_tdep (gdbarch);
+
+ amd_dbgapi_register_exists_t register_exists;
+ if (amd_dbgapi_wave_register_exists (wave_id, tdep->register_ids[regnum],
+ ®ister_exists)
+ != AMD_DBGAPI_STATUS_SUCCESS
+ || register_exists != AMD_DBGAPI_REGISTER_PRESENT)
+ return "";
+
+ return tdep->register_names[regnum].c_str ();
+}
+
+/* Return the internal register number for the DWARF register number DWARF_REG.
+
+ Return -1 if there's no internal register mapping to DWARF_REG. */
+
+static int
+amdgpu_dwarf_reg_to_regnum (struct gdbarch *gdbarch, int dwarf_reg)
+{
+ amdgpu_gdbarch_tdep *tdep = get_amdgpu_gdbarch_tdep (gdbarch);
+
+ if (dwarf_reg < tdep->dwarf_regnum_to_gdb_regnum.size ())
+ return tdep->dwarf_regnum_to_gdb_regnum[dwarf_reg];
+
+ return -1;
+}
+
+/* A hierarchy of classes to represent an amd-dbgapi register type. */
+
+struct amd_dbgapi_register_type
+{
+ enum class kind
+ {
+ INTEGER,
+ FLOAT,
+ DOUBLE,
+ VECTOR,
+ CODE_PTR,
+ FLAGS,
+ ENUM,
+ };
+
+ amd_dbgapi_register_type (kind kind, std::string lookup_name)
+ : m_kind (kind), m_lookup_name (std::move (lookup_name))
+ {}
+
+ virtual ~amd_dbgapi_register_type () = default;
+
+ /* Return the type's kind. */
+ kind kind () const
+ { return m_kind; }
+
+ /* Name to use for this type in the existing type map. */
+ const std::string &lookup_name () const
+ { return m_lookup_name; }
+
+private:
+ enum kind m_kind;
+ std::string m_lookup_name;
+};
+
+using amd_dbgapi_register_type_up = std::unique_ptr<amd_dbgapi_register_type>;
+
+struct amd_dbgapi_register_type_integer : public amd_dbgapi_register_type
+{
+ amd_dbgapi_register_type_integer (bool is_unsigned, unsigned int bit_size)
+ : amd_dbgapi_register_type
+ (kind::INTEGER,
+ string_printf ("%sint%d", is_unsigned ? "u" : "", bit_size)),
+ m_is_unsigned (is_unsigned),
+ m_bit_size (bit_size)
+ {}
+
+ bool is_unsigned () const
+ { return m_is_unsigned; }
+
+ unsigned int bit_size () const
+ { return m_bit_size; }
+
+private:
+ bool m_is_unsigned;
+ unsigned int m_bit_size;
+};
+
+struct amd_dbgapi_register_type_float : public amd_dbgapi_register_type
+{
+ amd_dbgapi_register_type_float ()
+ : amd_dbgapi_register_type (kind::FLOAT, "float")
+ {}
+};
+
+struct amd_dbgapi_register_type_double : public amd_dbgapi_register_type
+{
+ amd_dbgapi_register_type_double ()
+ : amd_dbgapi_register_type (kind::DOUBLE, "double")
+ {}
+};
+
+struct amd_dbgapi_register_type_vector : public amd_dbgapi_register_type
+{
+ amd_dbgapi_register_type_vector (const amd_dbgapi_register_type &element_type,
+ unsigned int count)
+ : amd_dbgapi_register_type (kind::VECTOR,
+ make_lookup_name (element_type, count)),
+ m_element_type (element_type),
+ m_count (count)
+ {}
+
+ const amd_dbgapi_register_type &element_type () const
+ { return m_element_type; }
+
+ unsigned int count () const
+ { return m_count; }
+
+ static std::string make_lookup_name
+ (const amd_dbgapi_register_type &element_type, unsigned int count)
+ {
+ return string_printf ("%s[%d]", element_type.lookup_name ().c_str (),
+ count);
+ }
+
+private:
+ const amd_dbgapi_register_type &m_element_type;
+ unsigned int m_count;
+};
+
+struct amd_dbgapi_register_type_code_ptr : public amd_dbgapi_register_type
+{
+ amd_dbgapi_register_type_code_ptr ()
+ : amd_dbgapi_register_type (kind::CODE_PTR, "void (*)()")
+ {}
+};
+
+struct amd_dbgapi_register_type_flags : public amd_dbgapi_register_type
+{
+ struct field
+ {
+ std::string name;
+ unsigned int bit_pos_start;
+ unsigned int bit_pos_end;
+ const amd_dbgapi_register_type *type;
+ };
+
+ using container_type = std::vector<field>;
+ using const_iterator_type = container_type::const_iterator;
+
+ amd_dbgapi_register_type_flags (unsigned int bit_size, gdb::string_view name)
+ : amd_dbgapi_register_type (kind::FLAGS,
+ make_lookup_name (bit_size, name)),
+ m_bit_size (bit_size),
+ m_name (std::move (name))
+ {}
+
+ unsigned int bit_size () const
+ { return m_bit_size; }
+
+ void add_field (std::string name, unsigned int bit_pos_start,
+ unsigned int bit_pos_end,
+ const amd_dbgapi_register_type *type)
+ {
+ m_fields.push_back (field {std::move (name), bit_pos_start,
+ bit_pos_end, type});
+ }
+
+ container_type::size_type size () const
+ { return m_fields.size (); }
+
+ const field &operator[] (container_type::size_type pos) const
+ { return m_fields[pos]; }
+
+ const_iterator_type begin () const
+ { return m_fields.begin (); }
+
+ const_iterator_type end () const
+ { return m_fields.end (); }
+
+ const std::string &name () const
+ { return m_name; }
+
+ static std::string make_lookup_name (int bits, gdb::string_view name)
+ {
+ std::string res = string_printf ("flags%d_t ", bits);
+ res.append (name.data (), name.size ());
+ return res;
+ }
+
+private:
+ unsigned int m_bit_size;
+ container_type m_fields;
+ std::string m_name;
+};
+
+using amd_dbgapi_register_type_flags_up
+ = std::unique_ptr<amd_dbgapi_register_type_flags>;
+
+struct amd_dbgapi_register_type_enum : public amd_dbgapi_register_type
+{
+ struct enumerator
+ {
+ std::string name;
+ ULONGEST value;
+ };
+
+ using container_type = std::vector<enumerator>;
+ using const_iterator_type = container_type::const_iterator;
+
+ amd_dbgapi_register_type_enum (gdb::string_view name)
+ : amd_dbgapi_register_type (kind::ENUM, make_lookup_name (name)),
+ m_name (name.data (), name.length ())
+ {}
+
+ void set_bit_size (int bit_size)
+ { m_bit_size = bit_size; }
+
+ unsigned int bit_size () const
+ { return m_bit_size; }
+
+ void add_enumerator (std::string name, ULONGEST value)
+ { m_enumerators.push_back (enumerator {std::move (name), value}); }
+
+ container_type::size_type size () const
+ { return m_enumerators.size (); }
+
+ const enumerator &operator[] (container_type::size_type pos) const
+ { return m_enumerators[pos]; }
+
+ const_iterator_type begin () const
+ { return m_enumerators.begin (); }
+
+ const_iterator_type end () const
+ { return m_enumerators.end (); }
+
+ const std::string &name () const
+ { return m_name; }
+
+ static std::string make_lookup_name (gdb::string_view name)
+ {
+ std::string res = "enum ";
+ res.append (name.data (), name.length ());
+ return res;
+ }
+
+private:
+ unsigned int m_bit_size = 32;
+ container_type m_enumerators;
+ std::string m_name;
+};
+
+using amd_dbgapi_register_type_enum_up
+ = std::unique_ptr<amd_dbgapi_register_type_enum>;
+
+/* Map type lookup names to types. */
+using amd_dbgapi_register_type_map
+ = std::unordered_map<std::string, amd_dbgapi_register_type_up>;
+
+/* Parse S as a ULONGEST, raise an error on overflow. */
+
+static ULONGEST
+try_strtoulst (gdb::string_view s)
+{
+ errno = 0;
+ ULONGEST value = strtoulst (s.data (), nullptr, 0);
+ if (errno != 0)
+ error (_("Failed to parse integer."));
+
+ return value;
+};
+
+/* Shared regex bits. */
+#define IDENTIFIER "[A-Za-z0-9_.]+"
+#define WS "[ \t]+"
+#define WSOPT "[ \t]*"
+
+static const amd_dbgapi_register_type &
+parse_amd_dbgapi_register_type (gdb::string_view type_name,
+ amd_dbgapi_register_type_map &type_map);
+
+
+/* parse_amd_dbgapi_register_type helper for enum types. */
+
+static void
+parse_amd_dbgapi_register_type_enum_fields
+ (amd_dbgapi_register_type_enum &enum_type, gdb::string_view fields)
+{
+ compiled_regex regex (/* name */
+ "^(" IDENTIFIER ")"
+ WSOPT "=" WSOPT
+ /* value */
+ "([0-9]+)"
+ WSOPT "(," WSOPT ")?",
+ REG_EXTENDED,
+ _("Error in AMDGPU enum register type regex"));
+ regmatch_t matches[4];
+
+ while (!fields.empty ())
+ {
+ int res = regex.exec (fields.data (), ARRAY_SIZE (matches), matches, 0);
+ if (res == REG_NOMATCH)
+ error (_("Failed to parse enum fields"));
+
+ auto sv_from_match = [fields] (const regmatch_t &m)
+ { return fields.substr (m.rm_so, m.rm_eo - m.rm_so); };
+
+ gdb::string_view name = sv_from_match (matches[1]);
+ gdb::string_view value_str = sv_from_match (matches[2]);
+ ULONGEST value = try_strtoulst (value_str);
+
+ if (value > std::numeric_limits<uint32_t>::max ())
+ enum_type.set_bit_size (64);
+
+ enum_type.add_enumerator (gdb::to_string (name), value);
+
+ fields = fields.substr (matches[0].rm_eo);
+ }
+}
+
+/* parse_amd_dbgapi_register_type helper for flags types. */
+
+static void
+parse_amd_dbgapi_register_type_flags_fields
+ (amd_dbgapi_register_type_flags &flags_type,
+ int bits, gdb::string_view name, gdb::string_view fields,
+ amd_dbgapi_register_type_map &type_map)
+{
+ gdb_assert (bits == 32 || bits == 64);
+
+ std::string regex_str
+ = string_printf (/* type */
+ "^(bool|uint%d_t|enum" WS IDENTIFIER WSOPT "(\\{[^}]*})?)"
+ WS
+ /* name */
+ "(" IDENTIFIER ")" WSOPT
+ /* bit position */
+ "@([0-9]+)(-[0-9]+)?" WSOPT ";" WSOPT,
+ bits);
+ compiled_regex regex (regex_str.c_str (), REG_EXTENDED,
+ _("Error in AMDGPU register type flags fields regex"));
+ regmatch_t matches[6];
+
+ while (!fields.empty ())
+ {
+ int res = regex.exec (fields.data (), ARRAY_SIZE (matches), matches, 0);
+ if (res == REG_NOMATCH)
+ error (_("Failed to parse flags type fields string"));
+
+ auto sv_from_match = [fields] (const regmatch_t &m)
+ { return fields.substr (m.rm_so, m.rm_eo - m.rm_so); };
+
+ gdb::string_view field_type_str = sv_from_match (matches[1]);
+ gdb::string_view field_name = sv_from_match (matches[3]);
+ gdb::string_view pos_begin_str = sv_from_match (matches[4]);
+ ULONGEST pos_begin = try_strtoulst (pos_begin_str);
+
+ if (field_type_str == "bool")
+ flags_type.add_field (gdb::to_string (field_name), pos_begin, pos_begin,
+ nullptr);
+ else
+ {
+ if (matches[5].rm_so == -1)
+ error (_("Missing end bit position"));
+
+ gdb::string_view pos_end_str = sv_from_match (matches[5]);
+ ULONGEST pos_end = try_strtoulst (pos_end_str.substr (1));
+ const amd_dbgapi_register_type &field_type
+ = parse_amd_dbgapi_register_type (field_type_str, type_map);
+ flags_type.add_field (gdb::to_string (field_name), pos_begin, pos_end,
+ &field_type);
+ }
+
+ fields = fields.substr (matches[0].rm_eo);
+ }
+}
+
+/* parse_amd_dbgapi_register_type helper for scalars. */
+
+static const amd_dbgapi_register_type &
+parse_amd_dbgapi_register_type_scalar (gdb::string_view name,
+ amd_dbgapi_register_type_map &type_map)
+{
+ std::string name_str = gdb::to_string (name);
+ auto it = type_map.find (name_str);
+ if (it != type_map.end ())
+ {
+ enum amd_dbgapi_register_type::kind kind = it->second->kind ();
+ if (kind != amd_dbgapi_register_type::kind::INTEGER
+ && kind != amd_dbgapi_register_type::kind::FLOAT
+ && kind != amd_dbgapi_register_type::kind::DOUBLE
+ && kind != amd_dbgapi_register_type::kind::CODE_PTR)
+ error (_("type mismatch"));
+
+ return *it->second;
+ }
+
+ amd_dbgapi_register_type_up type;
+ if (name == "int32_t")
+ type.reset (new amd_dbgapi_register_type_integer (false, 32));
+ else if (name == "uint32_t")
+ type.reset (new amd_dbgapi_register_type_integer (true, 32));
+ else if (name == "int64_t")
+ type.reset (new amd_dbgapi_register_type_integer (false, 64));
+ else if (name == "uint64_t")
+ type.reset (new amd_dbgapi_register_type_integer (true, 64));
+ else if (name == "float")
+ type.reset (new amd_dbgapi_register_type_float ());
+ else if (name == "double")
+ type.reset (new amd_dbgapi_register_type_double ());
+ else if (name == "void (*)()")
+ type.reset (new amd_dbgapi_register_type_code_ptr ());
+ else
+ error (_("unknown type %s"), name_str.c_str ());
+
+ auto insertion_pair = type_map.emplace (name, std::move (type));
+ return *insertion_pair.first->second;
+}
+
+/* Parse an amd-dbgapi register type string into an amd_dbgapi_register_type
+ object.
+
+ See the documentation of AMD_DBGAPI_REGISTER_INFO_TYPE in amd-dbgapi.h for
+ details about the format. */
+
+static const amd_dbgapi_register_type &
+parse_amd_dbgapi_register_type (gdb::string_view type_str,
+ amd_dbgapi_register_type_map &type_map)
+{
+ size_t pos_open_bracket = type_str.find_last_of ('[');
+ auto sv_from_match = [type_str] (const regmatch_t &m)
+ { return type_str.substr (m.rm_so, m.rm_eo - m.rm_so); };
+
+ if (pos_open_bracket != gdb::string_view::npos)
+ {
+ /* Vector types. */
+ gdb::string_view element_type_str
+ = type_str.substr (0, pos_open_bracket);
+ const amd_dbgapi_register_type &element_type
+ = parse_amd_dbgapi_register_type (element_type_str, type_map);
+
+ size_t pos_close_bracket = type_str.find_last_of (']');
+ gdb_assert (pos_close_bracket != gdb::string_view::npos);
+ gdb::string_view count_str_view
+ = type_str.substr (pos_open_bracket + 1,
+ pos_close_bracket - pos_open_bracket);
+ std::string count_str = gdb::to_string (count_str_view);
+ unsigned int count = std::stoul (count_str);
+
+ std::string lookup_name
+ = amd_dbgapi_register_type_vector::make_lookup_name (element_type, count);
+ auto existing_type_it = type_map.find (lookup_name);
+ if (existing_type_it != type_map.end ())
+ {
+ gdb_assert (existing_type_it->second->kind ()
+ == amd_dbgapi_register_type::kind::VECTOR);
+ return *existing_type_it->second;
+ }
+
+ amd_dbgapi_register_type_up type
+ (new amd_dbgapi_register_type_vector (element_type, count));
+ auto insertion_pair
+ = type_map.emplace (type->lookup_name (), std::move (type));
+ return *insertion_pair.first->second;
+ }
+
+ if (type_str.find ("flags32_t") == 0 || type_str.find ("flags64_t") == 0)
+ {
+ /* Split 'type_str' into 4 tokens: "(type) (name) ({ (fields) })". */
+ compiled_regex regex ("^(flags32_t|flags64_t)"
+ WS "(" IDENTIFIER ")" WSOPT
+ "(\\{" WSOPT "(.*)})?",
+ REG_EXTENDED,
+ _("Error in AMDGPU register type regex"));
+
+ regmatch_t matches[5];
+ int res = regex.exec (type_str.data (), ARRAY_SIZE (matches), matches, 0);
+ if (res == REG_NOMATCH)
+ error (_("Failed to parse flags type string"));
+
+ gdb::string_view flags_keyword = sv_from_match (matches[1]);
+ unsigned int bit_size = flags_keyword == "flags32_t" ? 32 : 64;
+ gdb::string_view name = sv_from_match (matches[2]);
+ std::string lookup_name
+ = amd_dbgapi_register_type_flags::make_lookup_name (bit_size, name);
+ auto existing_type_it = type_map.find (lookup_name);
+
+ if (matches[3].rm_so == -1)
+ {
+ /* No braces, lookup existing type. */
+ if (existing_type_it == type_map.end ())
+ error (_("reference to unknown type %s."),
+ gdb::to_string (name).c_str ());
+
+ if (existing_type_it->second->kind ()
+ != amd_dbgapi_register_type::kind::FLAGS)
+ error (_("type mismatch"));
+
+ return *existing_type_it->second;
+ }
+ else
+ {
+ /* With braces, it's a definition. */
+ if (existing_type_it != type_map.end ())
+ error (_("re-definition of type %s."),
+ gdb::to_string (name).c_str ());
+
+ amd_dbgapi_register_type_flags_up flags_type
+ (new amd_dbgapi_register_type_flags (bit_size, name));
+ gdb::string_view fields_without_braces = sv_from_match (matches[4]);
+
+ parse_amd_dbgapi_register_type_flags_fields
+ (*flags_type, bit_size, name, fields_without_braces, type_map);
+
+ auto insertion_pair
+ = type_map.emplace (flags_type->lookup_name (),
+ std::move (flags_type));
+ return *insertion_pair.first->second;
+ }
+ }
+
+ if (type_str.find ("enum") == 0)
+ {
+ compiled_regex regex ("^enum" WS "(" IDENTIFIER ")" WSOPT "(\\{" WSOPT "([^}]*)})?",
+ REG_EXTENDED,
+ _("Error in AMDGPU register type enum regex"));
+
+ /* Split 'type_name' into 3 tokens: "(name) ( { (fields) } )". */
+ regmatch_t matches[4];
+ int res = regex.exec (type_str.data (), ARRAY_SIZE (matches), matches, 0);
+ if (res == REG_NOMATCH)
+ error (_("Failed to parse flags type string"));
+
+ gdb::string_view name = sv_from_match (matches[1]);
+
+ std::string lookup_name
+ = amd_dbgapi_register_type_enum::make_lookup_name (name);
+ auto existing_type_it = type_map.find (lookup_name);
+
+ if (matches[2].rm_so == -1)
+ {
+ /* No braces, lookup existing type. */
+ if (existing_type_it == type_map.end ())
+ error (_("reference to unknown type %s"),
+ gdb::to_string (name).c_str ());
+
+ if (existing_type_it->second->kind ()
+ != amd_dbgapi_register_type::kind::ENUM)
+ error (_("type mismatch"));
+
+ return *existing_type_it->second;
+ }
+ else
+ {
+ /* With braces, it's a definition. */
+ if (existing_type_it != type_map.end ())
+ error (_("re-definition of type %s"),
+ gdb::to_string (name).c_str ());
+
+ amd_dbgapi_register_type_enum_up enum_type
+ (new amd_dbgapi_register_type_enum (name));
+ gdb::string_view fields_without_braces = sv_from_match (matches[3]);
+
+ parse_amd_dbgapi_register_type_enum_fields
+ (*enum_type, fields_without_braces);
+
+ auto insertion_pair
+ = type_map.emplace (enum_type->lookup_name (),
+ std::move (enum_type));
+ return *insertion_pair.first->second;
+ }
+ }
+
+ return parse_amd_dbgapi_register_type_scalar (type_str, type_map);
+}
+
+/* Convert an amd_dbgapi_register_type object to a GDB type. */
+
+static type *
+amd_dbgapi_register_type_to_gdb_type (const amd_dbgapi_register_type &type,
+ struct gdbarch *gdbarch)
+{
+ switch (type.kind ())
+ {
+ case amd_dbgapi_register_type::kind::INTEGER:
+ {
+ const auto &integer_type
+ = static_cast<const amd_dbgapi_register_type_integer &> (type);
+ switch (integer_type.bit_size ())
+ {
+ case 32:
+ if (integer_type.is_unsigned ())
+ return builtin_type (gdbarch)->builtin_uint32;
+ else
+ return builtin_type (gdbarch)->builtin_int32;
+
+ case 64:
+ if (integer_type.is_unsigned ())
+ return builtin_type (gdbarch)->builtin_uint64;
+ else
+ return builtin_type (gdbarch)->builtin_int64;
+
+ default:
+ gdb_assert_not_reached ("invalid bit size");
+ }
+ }
+
+ case amd_dbgapi_register_type::kind::VECTOR:
+ {
+ const auto &vector_type
+ = static_cast<const amd_dbgapi_register_type_vector &> (type);
+ struct type *element_type
+ = amd_dbgapi_register_type_to_gdb_type (vector_type.element_type (),
+ gdbarch);
+ return init_vector_type (element_type, vector_type.count ());
+ }
+
+ case amd_dbgapi_register_type::kind::FLOAT:
+ return builtin_type (gdbarch)->builtin_float;
+
+ case amd_dbgapi_register_type::kind::DOUBLE:
+ return builtin_type (gdbarch)->builtin_double;
+
+ case amd_dbgapi_register_type::kind::CODE_PTR:
+ return builtin_type (gdbarch)->builtin_func_ptr;
+
+ case amd_dbgapi_register_type::kind::FLAGS:
+ {
+ const auto &flags_type
+ = static_cast<const amd_dbgapi_register_type_flags &> (type);
+ struct type *gdb_type
+ = arch_flags_type (gdbarch, flags_type.name ().c_str (),
+ flags_type.bit_size ());
+
+ for (const auto &field : flags_type)
+ {
+ if (field.type == nullptr)
+ {
+ gdb_assert (field.bit_pos_start == field.bit_pos_end);
+ append_flags_type_flag (gdb_type, field.bit_pos_start,
+ field.name.c_str ());
+ }
+ else
+ {
+ struct type *field_type
+ = amd_dbgapi_register_type_to_gdb_type (*field.type, gdbarch);
+ gdb_assert (field_type != nullptr);
+ append_flags_type_field
+ (gdb_type, field.bit_pos_start,
+ field.bit_pos_end - field.bit_pos_start + 1,
+ field_type, field.name.c_str ());
+ }
+ }
+
+ return gdb_type;
+ }
+
+ case amd_dbgapi_register_type::kind::ENUM:
+ {
+ const auto &enum_type
+ = static_cast<const amd_dbgapi_register_type_enum &> (type);
+ struct type *gdb_type
+ = arch_type (gdbarch, TYPE_CODE_ENUM, enum_type.bit_size (),
+ enum_type.name ().c_str ());
+
+ gdb_type->set_num_fields (enum_type.size ());
+ gdb_type->set_fields
+ ((struct field *) TYPE_ZALLOC (gdb_type, (sizeof (struct field)
+ * enum_type.size ())));
+ gdb_type->set_is_unsigned (true);
+
+ for (size_t i = 0; i < enum_type.size (); ++i)
+ {
+ const auto &field = enum_type[i];
+ gdb_type->field (i).set_name (xstrdup (field.name.c_str ()));
+ gdb_type->field (i).set_loc_enumval (field.value);
+ }
+
+ return gdb_type;
+ }
+
+ default:
+ gdb_assert_not_reached ("unhandled amd_dbgapi_register_type kind");
+ }
+}
+
+static type *
+amdgpu_register_type (struct gdbarch *gdbarch, int regnum)
+{
+ amdgpu_gdbarch_tdep *tdep = get_amdgpu_gdbarch_tdep (gdbarch);
+
+ if (tdep->register_types[regnum] == nullptr)
+ {
+ /* This is done lazily (not at gdbarch initialization time), because it
+ requires access to builtin_type, which can't be used while the gdbarch
+ is not fully initialized. */
+ char *bytes;
+ amd_dbgapi_status_t status
+ = amd_dbgapi_register_get_info (tdep->register_ids[regnum],
+ AMD_DBGAPI_REGISTER_INFO_TYPE,
+ sizeof (bytes), &bytes);
+ if (status != AMD_DBGAPI_STATUS_SUCCESS)
+ error (_("Failed to get register type from amd-dbgapi"));
+
+ gdb::unique_xmalloc_ptr<char> bytes_holder (bytes);
+ amd_dbgapi_register_type_map type_map;
+ const amd_dbgapi_register_type ®ister_type
+ = parse_amd_dbgapi_register_type (bytes, type_map);
+ tdep->register_types[regnum]
+ = amd_dbgapi_register_type_to_gdb_type (register_type, gdbarch);
+ gdb_assert (tdep->register_types[regnum] != nullptr);
+ }
+
+ return tdep->register_types[regnum];
+}
+
+static int
+amdgpu_register_reggroup_p (struct gdbarch *gdbarch, int regnum,
+ const reggroup *group)
+{
+ amdgpu_gdbarch_tdep *tdep = get_amdgpu_gdbarch_tdep (gdbarch);
+
+ auto it = tdep->register_class_map.find (group->name ());
+ if (it == tdep->register_class_map.end ())
+ return group == all_reggroup;
+
+ amd_dbgapi_register_class_state_t state;
+ if (amd_dbgapi_register_is_in_register_class (it->second,
+ tdep->register_ids[regnum],
+ &state)
+ != AMD_DBGAPI_STATUS_SUCCESS)
+ return group == all_reggroup;
+
+ return (state == AMD_DBGAPI_REGISTER_CLASS_STATE_MEMBER
+ || group == all_reggroup);
+}
+
+static int
+amdgpu_breakpoint_kind_from_pc (struct gdbarch *gdbarch, CORE_ADDR *)
+{
+ return get_amdgpu_gdbarch_tdep (gdbarch)->breakpoint_instruction_size;
+}
+
+static const gdb_byte *
+amdgpu_sw_breakpoint_from_kind (struct gdbarch *gdbarch, int kind, int *size)
+{
+ *size = kind;
+ return get_amdgpu_gdbarch_tdep (gdbarch)->breakpoint_instruction_bytes.get ();
+}
+
+struct amdgpu_frame_cache
+{
+ CORE_ADDR base;
+ CORE_ADDR pc;
+};
+
+static amdgpu_frame_cache *
+amdgpu_frame_cache (frame_info_ptr this_frame, void **this_cache)
+{
+ if (*this_cache != nullptr)
+ return (struct amdgpu_frame_cache *) *this_cache;
+
+ struct amdgpu_frame_cache *cache
+ = FRAME_OBSTACK_ZALLOC (struct amdgpu_frame_cache);
+ (*this_cache) = cache;
+
+ cache->pc = get_frame_func (this_frame);
+ cache->base = 0;
+
+ return cache;
+}
+
+static void
+amdgpu_frame_this_id (frame_info_ptr this_frame, void **this_cache,
+ frame_id *this_id)
+{
+ struct amdgpu_frame_cache *cache
+ = amdgpu_frame_cache (this_frame, this_cache);
+
+ if (get_frame_type (this_frame) == INLINE_FRAME)
+ (*this_id) = frame_id_build (cache->base, cache->pc);
+ else
+ (*this_id) = outer_frame_id;
+
+ frame_debug_printf ("this_frame=%d, type=%d, this_id=%s",
+ frame_relative_level (this_frame),
+ get_frame_type (this_frame),
+ this_id->to_string ().c_str ());
+}
+
+static frame_id
+amdgpu_dummy_id (struct gdbarch *gdbarch, frame_info_ptr this_frame)
+{
+ return frame_id_build (0, get_frame_pc (this_frame));
+}
+
+static struct value *
+amdgpu_frame_prev_register (frame_info_ptr this_frame, void **this_cache,
+ int regnum)
+{
+ return frame_unwind_got_register (this_frame, regnum, regnum);
+}
+
+static const frame_unwind amdgpu_frame_unwi[...]
[diff truncated at 100000 bytes]
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2023-02-02 15:09 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2023-02-02 15:09 [binutils-gdb] gdb: initial support for ROCm platform (AMDGPU) debugging Simon Marchi
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).