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

[PATCH] 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.
---
 gdb/amd-dbgapi-target.c                     |  7 +++
 gdb/testsuite/gdb.rocm/interrupt-single.cpp | 43 +++++++++++++++++
 gdb/testsuite/gdb.rocm/interrupt-single.exp | 53 +++++++++++++++++++++
 3 files changed, 103 insertions(+)
 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/testsuite/gdb.rocm/interrupt-single.cpp b/gdb/testsuite/gdb.rocm/interrupt-single.cpp
new file mode 100644
index 00000000000..cbb78aed465
--- /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, 128>>> ();
+  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..fbe357abf1b
--- /dev/null
+++ b/gdb/testsuite/gdb.rocm/interrupt-single.exp
@@ -0,0 +1,53 @@
+# 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_main]} {
+	return
+    }
+
+    gdb_test "with breakpoint pending on -- tbreak loop" \
+	"breakpoint $::decimal \\(loop\\) pending."
+
+    gdb_test "continue" "breakpoint \[^\r\n\]+ loop.*"
+
+    # 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] gdb/amd-dbgapi-target: implement the 'interrupt' target op

[Lancelot SIX]

Hi Baris,

Thanks for looking into this.

On 11/05/2026 15:59, Tankut Baris Aktemur wrote:
> 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.
> ---
>   gdb/amd-dbgapi-target.c                     |  7 +++
>   gdb/testsuite/gdb.rocm/interrupt-single.cpp | 43 +++++++++++++++++
>   gdb/testsuite/gdb.rocm/interrupt-single.exp | 53 +++++++++++++++++++++
>   3 files changed, 103 insertions(+)
>   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/testsuite/gdb.rocm/interrupt-single.cpp b/gdb/testsuite/gdb.rocm/interrupt-single.cpp
> new file mode 100644
> index 00000000000..cbb78aed465
> --- /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 ();

We could as well drop the `loop` function and inline its loop here. 
That being said, it does not really hurt to have it either, so this is 
up to you.

> +}
> +
> +int
> +main ()
> +{
> +  /* Make sure that if anything goes wrong, the program eventually
> +     gets killed.  */
> +  gdb_watchdog (30);
> +
> +  kern<<<1, 128>>> ();

To keep the test minimal, could be <<<1, 1>>> (this will spawn 2 or 4 
waves depending on the generation), but again, this works as well.  I'll 
leave it up to you.

Otherwise, this looks good to me. Thanks.

Approved-by: Lancelot Six <lancelot.six@amd.com> (amdgpu)

> +  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..fbe357abf1b
> --- /dev/null
> +++ b/gdb/testsuite/gdb.rocm/interrupt-single.exp
> @@ -0,0 +1,53 @@
> +# 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_main]} {
> +	return
> +    }
> +
> +    gdb_test "with breakpoint pending on -- tbreak loop" \
> +	"breakpoint $::decimal \\(loop\\) pending."
> +
> +    gdb_test "continue" "breakpoint \[^\r\n\]+ loop.*"
> +
> +    # 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
> +	}
> +    }
> +}

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

[Aktemur, Baris]

AMD General

Hi Lancelot,

> Hi Baris,
>
> Thanks for looking into this.
>
> On 11/05/2026 15:59, Tankut Baris Aktemur wrote:
> > 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.
> > ---
> >   gdb/amd-dbgapi-target.c                     |  7 +++
> >   gdb/testsuite/gdb.rocm/interrupt-single.cpp | 43 +++++++++++++++++
> >   gdb/testsuite/gdb.rocm/interrupt-single.exp | 53 +++++++++++++++++++++
> >   3 files changed, 103 insertions(+)
> >   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/testsuite/gdb.rocm/interrupt-single.cpp
> b/gdb/testsuite/gdb.rocm/interrupt-single.cpp
> > new file mode 100644
> > index 00000000000..cbb78aed465
> > --- /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 ();
>
> We could as well drop the `loop` function and inline its loop here.
> That being said, it does not really hurt to have it either, so this is
> up to you.

I did it this way because it was easier to do "tbreak loop" instead of
using a line number.

But then it occurred to me that I could've used `runto "loop"` to further
simplify the code.  I also noticed that in the documentation there is a
restriction, for which the 'interrupt' command can be an alternative
mitigation.  I'd like to add this to the documentation, too.

Hence, with both changes, I'm going to send a v2 of the patch.

> > +}
> > +
> > +int
> > +main ()
> > +{
> > +  /* Make sure that if anything goes wrong, the program eventually
> > +     gets killed.  */
> > +  gdb_watchdog (30);
> > +
> > +  kern<<<1, 128>>> ();
>
> To keep the test minimal, could be <<<1, 1>>> (this will spawn 2 or 4
> waves depending on the generation), but again, this works as well.  I'll
> leave it up to you.

Done.

> Otherwise, this looks good to me. Thanks.
>
> Approved-by: Lancelot Six <lancelot.six@amd.com> (amdgpu)

Thank you for the review.  I kindly ask that you take a look at v2, too.

Regards,
-Baris