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

https://sourceware.org/git/gitweb.cgi?p=binutils-gdb.git;h=3c5c7b7b4bfbc4c89a09678a2cbb43482a15aa6e

commit 3c5c7b7b4bfbc4c89a09678a2cbb43482a15aa6e
Author: Lancelot SIX <[email protected]>
Date:   Wed Aug 5 22:45:37 2026 +0100

    gdb/amdgpu: Handle SIGABRT with a higher priority than SIGTRAP
    
    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 expect 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
    Reviewed-By: Tankut Baris Aktemur <[email protected]>

Diff:
---
 gdb/amd-dbgapi-target.c               |  4 +-
 gdb/testsuite/gdb.rocm/step-abort.cpp | 32 ++++++++++++++++
 gdb/testsuite/gdb.rocm/step-abort.exp | 71 +++++++++++++++++++++++++++++++++++
 3 files changed, 105 insertions(+), 2 deletions(-)

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..56bcdfc8534
--- /dev/null
+++ b/gdb/testsuite/gdb.rocm/step-abort.exp
@@ -0,0 +1,71 @@
+# 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 $::testfile
+
+    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