[PATCH] gdb/amdgpu: Handle SIGABRT with a higher priority than SIGTRAP

[PATCH] gdb/amdgpu: Handle SIGABRT with a higher priority than SIGTRAP

[Lancelot SIX]

On the AMDGPU target, waves (known as threads by GDB) can report
multiple events at the same time.  However, the amd-dbgapi-target can
only report one target_waitstatus to the core of GDB.  This means that
when multiple exceptions are reported at once, the target needs to
choose which one is the most important.

In the current implementation, if we single step the instruction which
should cause a STOP_REASON_ABORT, the target only reports the single
step (GDB_SIGNAL_TRAP), missing the abort signal (GDB_SIGNAL_ABRT).
However, when single stepping an abort, we expert SIGABRT to be shown to
the user.

This patch proposes to change the priority in the target so
STOP_REASON_ASSERT_TRAP takes priority over STOP_REASON_SINGLE_STEP and
other debugger related traps such as watchpoint.

Add a testcase which have GDB single step a simple shader until it calls
abort ().  Before this patch, we had:

    (gdb) x/3i $pc
    => 0x7ffff7fa9600 <_Z4kernv>:   s_sleep 8
       0x7ffff7fa9604 <_Z4kernv+4>: s_trap 2    # The abort instruction
       0x7ffff7fa9608:      v_illegal
    (gdb) si
    0x00007ffff7fa9604 in kern() () from file:///.../step-abort#offset=8192&size=3296
    (gdb) si
    0x00007ffff7fa9608 in ?? ()
    (gdb) si

    Thread 5 "kern" received signal SIGILL, Illegal instruction.
    0x00007ffff7fa960c in ?? ()

GDB would single step over the s_trap 2 instruction, but silently hide
the SIGABRT, trying to execute past the end of the shader.  With this
patch, GDB correctly recognises the abort:

    (gdb) si
    0x00007ffff7fa9604 in kern() ()
       from file:///.../step-abort#offset=8192&size=3296
    (gdb) si

    Thread 5 "kern" received signal SIGABRT, Aborted.
    0x00007ffff7fa9608 in ?? ()

Since the SIGABRT is now correctly reported to GDB, the next continue
will be able to resume the thread with the appropriate signal, notifying
the runtime that the queue where the shader was running is now in the
error state.

Tested on x86_64-linux + AMDGPU gfx1031.

Change-Id: I0223769816dfe08b92d7401c56b99ec6e46369bd
---
 gdb/amd-dbgapi-target.c               |  4 +-
 gdb/testsuite/gdb.rocm/step-abort.cpp | 32 ++++++++++++
 gdb/testsuite/gdb.rocm/step-abort.exp | 72 +++++++++++++++++++++++++++
 3 files changed, 106 insertions(+), 2 deletions(-)
 create mode 100644 gdb/testsuite/gdb.rocm/step-abort.cpp
 create mode 100644 gdb/testsuite/gdb.rocm/step-abort.exp

diff --git a/gdb/amd-dbgapi-target.c b/gdb/amd-dbgapi-target.c
index b4ca1506906..9c6cc99a43f 100644
--- a/gdb/amd-dbgapi-target.c
+++ b/gdb/amd-dbgapi-target.c
@@ -1505,6 +1505,8 @@ process_one_event (amd_dbgapi_inferior_info &info,
 			| AMD_DBGAPI_WAVE_STOP_REASON_FP_INVALID_OPERATION
 			| AMD_DBGAPI_WAVE_STOP_REASON_INT_DIVIDE_BY_0))
 	      ws.set_stopped (GDB_SIGNAL_FPE);
+	    else if (stop_reason & AMD_DBGAPI_WAVE_STOP_REASON_ASSERT_TRAP)
+	      ws.set_stopped (GDB_SIGNAL_ABRT);
 	    else if (stop_reason
 		     & (AMD_DBGAPI_WAVE_STOP_REASON_BREAKPOINT
 			| AMD_DBGAPI_WAVE_STOP_REASON_WATCHPOINT
@@ -1512,8 +1514,6 @@ process_one_event (amd_dbgapi_inferior_info &info,
 			| AMD_DBGAPI_WAVE_STOP_REASON_DEBUG_TRAP
 			| AMD_DBGAPI_WAVE_STOP_REASON_TRAP))
 	      ws.set_stopped (GDB_SIGNAL_TRAP);
-	    else if (stop_reason & AMD_DBGAPI_WAVE_STOP_REASON_ASSERT_TRAP)
-	      ws.set_stopped (GDB_SIGNAL_ABRT);
 	    else
 	      ws.set_stopped (GDB_SIGNAL_0);
 
diff --git a/gdb/testsuite/gdb.rocm/step-abort.cpp b/gdb/testsuite/gdb.rocm/step-abort.cpp
new file mode 100644
index 00000000000..560697fde19
--- /dev/null
+++ b/gdb/testsuite/gdb.rocm/step-abort.cpp
@@ -0,0 +1,32 @@
+/* This testcase is part of GDB, the GNU debugger.
+
+   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/>.  */
+
+#include "hip/hip_runtime.h"
+
+__global__ void
+kern ()
+{
+  __builtin_amdgcn_s_sleep (8);
+  __builtin_abort ();
+}
+
+int
+main ()
+{
+  kern<<<1, 1>>> ();
+  return hipDeviceSynchronize () != hipSuccess;
+}
diff --git a/gdb/testsuite/gdb.rocm/step-abort.exp b/gdb/testsuite/gdb.rocm/step-abort.exp
new file mode 100644
index 00000000000..a385bc435e4
--- /dev/null
+++ b/gdb/testsuite/gdb.rocm/step-abort.exp
@@ -0,0 +1,72 @@
+# 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/>.
+
+# This test ensures that we receive SIGABRT when we step over an abort
+# instruction.
+
+load_lib rocm.exp
+
+standard_testfile .cpp
+
+require allow_hipcc_tests
+
+# We want to have a small kernel as we are going to single step all the way
+# to our abort instruction (s_trap 2).  Using -O1 allows the compiler to inline
+# the sleep and abort instructions.
+if {[build_executable "failed to prepare" $testfile $srcfile {hip additional_flags=-O1}]} {
+    return
+}
+
+proc do_test {} {
+    clean_restart
+    gdb_load $::binfile
+
+    with_rocm_gpu_lock {
+	if {![runto_main]} {
+	    return
+	}
+
+	gdb_test "with breakpoint pending on -- break kern" \
+	    "Breakpoint $::decimal \\(kern\\) pending."
+
+	gdb_test "continue" \
+	    "Thread $::decimal hit Breakpoint $::decimal.* kern.*"
+
+	set remaining_steps 60
+	gdb_test_multiple "si" "step until SIGABRT" {
+	    -re -wrap ".*SIGABRT.*" {
+		pass $gdb_test_name
+	    }
+	    -re -wrap ".*" {
+		incr remaining_steps -1
+		verbose -log "remaining steps: $remaining_steps"
+		if {$remaining_steps == 0} {
+		    fail $gdb_test_name
+		} else {
+		    send_gdb "si\n"
+		    exp_continue
+		}
+	    }
+	}
+
+	# We have received the SIGABRT.  If we continue from here, the
+	# exception is passed to the inferior, i.e. GDB forwards it to the
+	# ROCr runtime, which can detect the shader error.
+	gdb_test "continue" "Queue error: HSA_STATUS_ERROR_EXCEPTION.*" \
+	    "send exception to the runtime"
+    }
+}
+
+do_test

base-commit: 9790ec8b5538fdf7d61474c48c98bd9dbe0b4d31
-- 
2.43.0

RE: [PATCH] gdb/amdgpu: Handle SIGABRT with a higher priority than SIGTRAP

[Aktemur, Baris]

AMD General

Hi Lancelot,

On Thursday, August 6, 2026 12:24 AM, Six, Lancelot wrote:
> On the AMDGPU target, waves (known as threads by GDB) can report
> multiple events at the same time.  However, the amd-dbgapi-target can
> only report one target_waitstatus to the core of GDB.  This means that
> when multiple exceptions are reported at once, the target needs to
> choose which one is the most important.
>
> In the current implementation, if we single step the instruction which
> should cause a STOP_REASON_ABORT, the target only reports the single
> step (GDB_SIGNAL_TRAP), missing the abort signal (GDB_SIGNAL_ABRT).
> However, when single stepping an abort, we expert SIGABRT to be shown to

Typo: expert -> expect

> the user.
>
> This patch proposes to change the priority in the target so
> STOP_REASON_ASSERT_TRAP takes priority over STOP_REASON_SINGLE_STEP and
> other debugger related traps such as watchpoint.
>
> Add a testcase which have GDB single step a simple shader until it calls
> abort ().  Before this patch, we had:
>
>     (gdb) x/3i $pc
>     => 0x7ffff7fa9600 <_Z4kernv>:   s_sleep 8
>        0x7ffff7fa9604 <_Z4kernv+4>: s_trap 2    # The abort instruction
>        0x7ffff7fa9608:      v_illegal
>     (gdb) si
>     0x00007ffff7fa9604 in kern() () from file:///.../step-abort#offset=8192&size=3296
>     (gdb) si
>     0x00007ffff7fa9608 in ?? ()
>     (gdb) si
>
>     Thread 5 "kern" received signal SIGILL, Illegal instruction.
>     0x00007ffff7fa960c in ?? ()
>
> GDB would single step over the s_trap 2 instruction, but silently hide
> the SIGABRT, trying to execute past the end of the shader.  With this
> patch, GDB correctly recognises the abort:
>
>     (gdb) si
>     0x00007ffff7fa9604 in kern() ()
>        from file:///.../step-abort#offset=8192&size=3296
>     (gdb) si
>
>     Thread 5 "kern" received signal SIGABRT, Aborted.
>     0x00007ffff7fa9608 in ?? ()
>
> Since the SIGABRT is now correctly reported to GDB, the next continue
> will be able to resume the thread with the appropriate signal, notifying
> the runtime that the queue where the shader was running is now in the
> error state.
>
> Tested on x86_64-linux + AMDGPU gfx1031.

Looks good to me.  Thanks.
Reviewed-By: Tankut Baris Aktemur <TankutBaris.Aktemur@amd.com>

-Baris

Re: [PATCH] gdb/amdgpu: Handle SIGABRT with a higher priority than SIGTRAP

[Lancelot SIX]

> 
> Looks good to me.  Thanks.
> Reviewed-By: Tankut Baris Aktemur <TankutBaris.Aktemur@amd.com>
> 
> -Baris
> 

Thanks.

Suggestions applied, and pushed as 3c5c7b7b4bfbc4c89a09678a2cbb43482a15aa6e

Best,
Lancelot.

RE: [PATCH] gdb/amdgpu: Handle SIGABRT with a higher priority than SIGTRAP

[Aktemur, Baris]

AMD General

> diff --git a/gdb/testsuite/gdb.rocm/step-abort.exp b/gdb/testsuite/gdb.rocm/step-
> abort.exp
> new file mode 100644
> index 00000000000..a385bc435e4
> --- /dev/null
> +++ b/gdb/testsuite/gdb.rocm/step-abort.exp
> @@ -0,0 +1,72 @@
> +# 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/>.
> +
> +# This test ensures that we receive SIGABRT when we step over an abort
> +# instruction.
> +
> +load_lib rocm.exp
> +
> +standard_testfile .cpp
> +
> +require allow_hipcc_tests
> +
> +# We want to have a small kernel as we are going to single step all the way
> +# to our abort instruction (s_trap 2).  Using -O1 allows the compiler to inline
> +# the sleep and abort instructions.
> +if {[build_executable "failed to prepare" $testfile $srcfile {hip additional_flags=-
> O1}]} {
> +    return
> +}
> +
> +proc do_test {} {
> +    clean_restart
> +    gdb_load $::binfile

I forgot to mention a nit.  This can be a single line:

  clean_restart $::testfile

Regards,
-Baris