[PATCH v2] gdb/amd-dbgapi-target: implement the 'interrupt' target op

[PATCH v2] gdb/amd-dbgapi-target: implement the 'interrupt' target op

[Tankut Baris Aktemur]

In all-stop mode, an asynchronously-resumed AMD GPU thread cannot be
stopped using the "interrupt" command, because in that case the
interrupt request is simply passed to the target beneath (i.e. the
native CPU target), which does nothing since all CPU threads are
already stopped.

Implement the 'interrupt' method of amd_dbgapi_target by using the
'stop' method.  Also include a test.

Approved-by: Lancelot Six <lancelot.six@amd.com> (amdgpu)
---
 gdb/amd-dbgapi-target.c                     |  7 +++
 gdb/doc/gdb.texinfo                         |  4 +-
 gdb/testsuite/gdb.rocm/interrupt-single.cpp | 43 ++++++++++++++++++
 gdb/testsuite/gdb.rocm/interrupt-single.exp | 48 +++++++++++++++++++++
 4 files changed, 101 insertions(+), 1 deletion(-)
 create mode 100644 gdb/testsuite/gdb.rocm/interrupt-single.cpp
 create mode 100644 gdb/testsuite/gdb.rocm/interrupt-single.exp

diff --git a/gdb/amd-dbgapi-target.c b/gdb/amd-dbgapi-target.c
index 421ec8599ed..fe67484db12 100644
--- a/gdb/amd-dbgapi-target.c
+++ b/gdb/amd-dbgapi-target.c
@@ -299,6 +299,7 @@ struct amd_dbgapi_target final : public target_ops
   void resume (ptid_t, int, enum gdb_signal) override;
   void commit_resumed () override;
   void stop (ptid_t ptid) override;
+  void interrupt () override;
 
   void fetch_registers (struct regcache *, int) override;
   void store_registers (struct regcache *, int) override;
@@ -1165,6 +1166,12 @@ amd_dbgapi_target::stop (ptid_t ptid)
 	stop_one_thread (&thread);
 }
 
+void
+amd_dbgapi_target::interrupt ()
+{
+  stop (minus_one_ptid);
+}
+
 /* Callback for our async event handler.  */
 
 static void
diff --git a/gdb/doc/gdb.texinfo b/gdb/doc/gdb.texinfo
index f6b3b114278..d4d8a20b5bb 100644
--- a/gdb/doc/gdb.texinfo
+++ b/gdb/doc/gdb.texinfo
@@ -27872,7 +27872,9 @@ If no CPU thread is running, then @samp{Ctrl-C} is not able to stop
 @acronym{AMD GPU} threads.  This can happen for example if you enable
 @code{scheduler-locking} after the whole program stopped, and then resume an
 @acronym{AMD GPU} thread.  The only way to unblock the situation is to kill the
-@value{GDBN} process.
+@value{GDBN} process.  Alternatively you can resume the @acronym{AMD GPU}
+thread in the background using @code{continue&} (@pxref{Background Execution})
+and then use @code{interrupt} to stop the thread.
 
 @anchor{AMD GPU Attaching Restrictions}
 @item
diff --git a/gdb/testsuite/gdb.rocm/interrupt-single.cpp b/gdb/testsuite/gdb.rocm/interrupt-single.cpp
new file mode 100644
index 00000000000..fc8d2cca697
--- /dev/null
+++ b/gdb/testsuite/gdb.rocm/interrupt-single.cpp
@@ -0,0 +1,43 @@
+/* Copyright 2026 Free Software Foundation, Inc.
+
+   This file is part of GDB.
+
+   This program is free software; you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation; either version 3 of the License, or
+   (at your option) any later version.
+
+   This program is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with this program.  If not, see <http://www.gnu.org/licenses/>.  */
+
+#include <hip/hip_runtime.h>
+#include "gdb_watchdog.h"
+
+__device__ void
+loop ()
+{
+  while (true)
+    __builtin_amdgcn_s_sleep (8);
+}
+
+__global__ void
+kern ()
+{
+  loop ();
+}
+
+int
+main ()
+{
+  /* Make sure that if anything goes wrong, the program eventually
+     gets killed.  */
+  gdb_watchdog (30);
+
+  kern<<<1, 1>>> ();
+  return hipDeviceSynchronize () != hipSuccess;
+}
diff --git a/gdb/testsuite/gdb.rocm/interrupt-single.exp b/gdb/testsuite/gdb.rocm/interrupt-single.exp
new file mode 100644
index 00000000000..23ae02a0c74
--- /dev/null
+++ b/gdb/testsuite/gdb.rocm/interrupt-single.exp
@@ -0,0 +1,48 @@
+# Copyright 2026 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/>.
+
+# Test that a running thread can be stopped with the 'interrupt'
+# command.  This is done in all-stop mode to ensure the 'interrupt'
+# target op is used.
+
+load_lib rocm.exp
+
+require allow_hipcc_tests
+
+standard_testfile .cpp
+
+if {[prepare_for_testing "failed to prepare" $testfile $srcfile {debug hip}]} {
+    return
+}
+
+with_rocm_gpu_lock {
+    if {![runto "loop" {allow-pending}]} {
+	return
+    }
+
+    # Resume a single GPU thread asynchronously to be able to use
+    # the interrupt command.
+    gdb_test_no_output "set scheduler-locking on"
+    gdb_test "continue &" "Continuing\." "continue async"
+
+    gdb_test_multiple "interrupt" "" {
+	-re "Thread $::decimal stopped" {
+	    pass $gdb_test_name
+	}
+	-re "$::gdb_prompt" {
+	    exp_continue
+	}
+    }
+}
-- 
2.34.1

Re: [PATCH v2] gdb/amd-dbgapi-target: implement the 'interrupt' target op

[Eli Zaretskii]

> From: Tankut Baris Aktemur <tankutbaris.aktemur@amd.com>
> CC: <lancelot.six@amd.com>, <eliz@gnu.org>
> Date: Tue, 12 May 2026 12:33:18 -0500
> 
> In all-stop mode, an asynchronously-resumed AMD GPU thread cannot be
> stopped using the "interrupt" command, because in that case the
> interrupt request is simply passed to the target beneath (i.e. the
> native CPU target), which does nothing since all CPU threads are
> already stopped.
> 
> Implement the 'interrupt' method of amd_dbgapi_target by using the
> 'stop' method.  Also include a test.
> 
> Approved-by: Lancelot Six <lancelot.six@amd.com> (amdgpu)
> ---
>  gdb/amd-dbgapi-target.c                     |  7 +++
>  gdb/doc/gdb.texinfo                         |  4 +-
>  gdb/testsuite/gdb.rocm/interrupt-single.cpp | 43 ++++++++++++++++++
>  gdb/testsuite/gdb.rocm/interrupt-single.exp | 48 +++++++++++++++++++++
>  4 files changed, 101 insertions(+), 1 deletion(-)
>  create mode 100644 gdb/testsuite/gdb.rocm/interrupt-single.cpp
>  create mode 100644 gdb/testsuite/gdb.rocm/interrupt-single.exp

Thanks.

> +@value{GDBN} process.  Alternatively you can resume the @acronym{AMD GPU}
                                                           ^^^^^^^^^^^^^^^^^
This should be 2 separate @acronym's.

> +thread in the background using @code{continue&} (@pxref{Background Execution})
> +and then use @code{interrupt} to stop the thread.

The documentation part is okay with this nit fixed.

Reviewed-By: Eli Zaretskii <eliz@gnu.org>

Re: [PATCH v2] gdb/amd-dbgapi-target: implement the 'interrupt' target op

[Pedro Alves]

Hi!

On 2026-05-12 18:33, Tankut Baris Aktemur wrote:

>  
> +void
> +amd_dbgapi_target::interrupt ()
> +{
> +  stop (minus_one_ptid);
> +}

Two issues here:

#1 - Note that this changes the behavior of the CPU target.  

We push the dbgapi target on the stack even before GPU debugging is active.  

Before the patch, "interrupt" stops with SIGINT.

After the patch, "interrupt" stops with SIGINT vs "stopped" depending on whether
dbgapi is built into GDB and active, or not.

Ultimately, I want to make the linux target interrupt with "stopped" too, and my
larger "ctrl-c rework" series does that:
  https://inbox.sourceware.org/gdb-patches/20210603190243.2609886-17-pedro@palves.net/

but maybe we can have a simpler localized patch in linux-nat.c meanwhile, similar to
what I did for Windows recently, here:
  https://inbox.sourceware.org/gdb-patches/20260505122434.1444507-1-pedro@palves.net/


#2 - That is going to iterate over all threads and queue a stop for every thread.  That is
incorrect in all-stop mode.  The assumption is that there is only one single event queued,
and then infrun stops all threads when it handles the event.

E.g., if I try without scheduler-locking, I see:

(gdb) c&
Continuing.
(gdb) info threads
  Id   Target Id                                             Frame 
  1    Thread 0x7ffff62a8540 (LWP 3244715) "interrupt-singl" (running)
  2    Thread 0x7fffe8dff6c0 (LWP 3244828) "interrupt-singl" (running)
  3    Thread 0x7fffe3fff6c0 (LWP 3244829) "interrupt-singl" (running)
  6    Thread 0x7ffff5adf6c0 (LWP 3244832) "interrupt-singl" (running)
* 7    AMDGPU Wave 1:2:1:1 (0,0,0)/0 "kern"                  (running)
(gdb) interrupt 
(gdb) ❌️ wave_stop for wave_1 failed (The wave has an outstanding stop request)
(gdb) ❌️ wave_stop for wave_1 failed (The wave has an outstanding stop request)
(gdb) ❌️ wave_stop for wave_1 failed (The wave has an outstanding stop request)
(gdb) ❌️ wave_stop for wave_1 failed (The wave has an outstanding stop request)
(gdb) 
Thread 7 "kern" stopped.
0x00007ffff7fa7d60 in loop () at /home/pedro/rocm/gdb/build/gdb/testsuite/../../../src/gdb/testsuite/gdb.rocm/interrupt-single.cpp:25
25          __builtin_amdgcn_s_sleep (8);


It's probably also possible to trigger badness in the CPU side.

Your test should definitely cover this, not just schedlock on.

Pedro Alves

RE: [PATCH v2] gdb/amd-dbgapi-target: implement the 'interrupt' target op

[Aktemur, Baris]

AMD General

Hi Eli,

> > +@value{GDBN} process.  Alternatively you can resume the @acronym{AMD GPU}
>                                                            ^^^^^^^^^^^^^^^^^
> This should be 2 separate @acronym's.

You mean it should be "@acronym{AMD} @acronym{GPU}"?

There are several occurrences of "@acronym{AMD ROCm} and @acronym{AMD GPU}"
in the document.  I'd send a patch to convert them to "@acronym{AMD}
@acronym{ROCm}" and "@acronym{AMD} @acronym{GPU}", respectively.  OK with that?

Regards,
-Baris

Re: [PATCH v2] gdb/amd-dbgapi-target: implement the 'interrupt' target op

[Eli Zaretskii]

> From: "Aktemur, Baris" <TankutBaris.Aktemur@amd.com>
> CC: "gdb-patches@sourceware.org" <gdb-patches@sourceware.org>, "Six, Lancelot"
> 	<Lancelot.Six@amd.com>
> Date: Wed, 13 May 2026 05:38:53 +0000
> 
> > > +@value{GDBN} process.  Alternatively you can resume the @acronym{AMD GPU}
> >                                                            ^^^^^^^^^^^^^^^^^
> > This should be 2 separate @acronym's.
> 
> You mean it should be "@acronym{AMD} @acronym{GPU}"?

Yes.

> There are several occurrences of "@acronym{AMD ROCm} and @acronym{AMD GPU}"
> in the document.  I'd send a patch to convert them to "@acronym{AMD}
> @acronym{ROCm}" and "@acronym{AMD} @acronym{GPU}", respectively.  OK with that?

Yes, please.