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

[Tankut Baris Aktemur <tankutbaris.aktemur@amd.com>]

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