gdb/amd-dbgapi: add displaced stepping support

Implement the target_ops displaced stepping methods to add displaced
stepping support when debugging AMD GPU programs.  The knowledge of how
to prepare and finish displaced steps is provided by the amd-dbgapi
library, so the code here is relatively straightforward.  No need to
parse instructions or handle fixups, that is done by the lib  We just
need to remember, for each thread doing a displaced step, the displaced
stepping id given by the library.

Add a test to exercise the new functionality.  The compiler generates
DWARF that GDB doesn't understand yet [1], so trying to step over a
breakpoint with DWARF present gives:

    (gdb) si
    Unhandled dwarf expression opcode 0xe9

The test purposefully builds the binary without DWARF info to circumvent
this.

[1] https://llvm.org/docs/AMDGPUDwarfExtensionsForHeterogeneousDebugging.html

Change-Id: I53f459221a42d4b02a6041eadb8cf554500e2162
Approved-By: Lancelot Six <lancelot.six@amd.com> (amdgpu)

1 files changed, 48 insertions, 0 deletions

diff --git a/gdb/testsuite/gdb.rocm/displaced-stepping.cpp b/gdb/testsuite/gdb.rocm/displaced-stepping.cpp
new file mode 100644
index 0000000..b94e8ab
--- /dev/null
+++ b/gdb/testsuite/gdb.rocm/displaced-stepping.cpp
@@ -0,0 +1,48 @@
+/* This testcase is part of GDB, the GNU debugger.
+
+   Copyright 2022-2024 Free Software Foundation, Inc.
+
+   This program is free software; you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation; either version 3 of the License, or
+   (at your option) any later version.
+
+   This program is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with this program.  If not, see <http://www.gnu.org/licenses/>.  */
+
+#include "hip/hip_runtime.h"
+#include <cassert>
+
+__global__ void
+do_an_addition (int a, int b, int *out)
+{
+  *out = a + b;
+}
+
+int
+main ()
+{
+  int *result_ptr, result;
+
+  /* Allocate memory for the device to write the result to.  */
+  hipError_t error = hipMalloc (&result_ptr, sizeof (int));
+  assert (error == hipSuccess);
+
+  /* Run `do_an_addition` on one workgroup containing one work item.  */
+  do_an_addition<<<dim3(1), dim3(1), 0, 0>>> (1, 2, result_ptr);
+
+  /* Copy result from device to host.  Note that this acts as a synchronization
+     point, waiting for the kernel dispatch to complete.  */
+  error = hipMemcpyDtoH (&result, result_ptr, sizeof (int));
+  assert (error == hipSuccess);
+
+  printf ("result is %d\n", result);
+  assert (result == 3);
+
+  return 0;
+}