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

["Aktemur, Baris" <TankutBaris.Aktemur@amd.com>]

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