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

[Lancelot SIX <Lancelot.Six@amd.com>]

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
> +	}
> +    }
> +}