From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from simark.ca by simark.ca with LMTP id EUWiNE64c2rqEQsAWB0awg (envelope-from ) for ; Wed, 05 Aug 2026 18:25:18 -0400 Authentication-Results: simark.ca; dkim=pass (1024-bit key; unprotected) header.d=amd.com header.i=@amd.com header.a=rsa-sha256 header.s=selector1 header.b=2GpRPVgc; dkim-atps=neutral Received: by simark.ca (Postfix, from userid 112) id C957E1E166; Wed, 05 Aug 2026 18:25:18 -0400 (EDT) X-Spam-Checker-Version: SpamAssassin 4.0.1 (2024-03-25) on simark.ca X-Spam-Level: X-Spam-Status: No, score=-6.4 required=5.0 tests=ARC_SIGNED,ARC_VALID,BAYES_00, DKIMWL_WL_HIGH,DKIM_SIGNED,DKIM_VALID,DKIM_VALID_AU,MAILING_LIST_MULTI, RCVD_IN_DNSWL_MED autolearn=ham autolearn_force=no version=4.0.1 Received: from vm01.sourceware.org (vm01.sourceware.org [IPv6:2620:52:6:3111::32]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange x25519 server-signature ECDSA (prime256v1) server-digest SHA256) (No client certificate requested) by simark.ca (Postfix) with ESMTPS id 6274D1E033 for ; Wed, 05 Aug 2026 18:25:17 -0400 (EDT) Received: from vm01.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 8A6084BA2E28 for ; Wed, 5 Aug 2026 22:25:16 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 8A6084BA2E28 Authentication-Results: sourceware.org; dkim=pass (1024-bit key, unprotected) header.d=amd.com header.i=@amd.com header.a=rsa-sha256 header.s=selector1 header.b=2GpRPVgc Received: from SJ2PR03CU001.outbound.protection.outlook.com (mail-westusazon11012070.outbound.protection.outlook.com [52.101.43.70]) by sourceware.org (Postfix) with ESMTPS id 77B4B4BAE7CD for ; Wed, 5 Aug 2026 22:24:49 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 77B4B4BAE7CD Authentication-Results: sourceware.org; dmarc=pass (p=quarantine dis=none) header.from=amd.com Authentication-Results: sourceware.org; spf=fail smtp.mailfrom=amd.com ARC-Filter: OpenARC Filter v1.0.0 sourceware.org 77B4B4BAE7CD Authentication-Results: sourceware.org; arc=pass smtp.remote-ip=52.101.43.70 ARC-Seal: i=2; a=rsa-sha256; d=sourceware.org; s=key; t=1785968690; cv=pass; b=Szp7tXx8yL4RHadsseKkUe1tR5zEtyMlt6UhPpNj+d11iMe25fl8bslyyQ2gl7pt0Z+FxBZd899gcfOnEIZIavjnF8YNMmMd4m8Ws3teVlfQHjCz0v9LtwNhL0PQLWVrwyB1hrv//7PV12owEpISgbpT8IMPgmcExZelwmOZOgE= ARC-Message-Signature: i=2; a=rsa-sha256; d=sourceware.org; s=key; t=1785968690; c=relaxed/simple; bh=MPPtDOL70kGwMq3HsqwfKzrlCyamZLHpa0cugY9TrAI=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=eo91xeA7QBHcBqJrs0yqJVAhMV13bHCC1GsPqZxH6jDolXaw2lO1315YZt5xuDIkxvAyxdZkPkmtIW4fNF34gObagzqoiARkPJmpVjt+YEcl+OFpRffE8FWN/szB3bQcMykWSsK8t39WTRPlUpDuT5zww69HdlAoqw/YGfYpSis= ARC-Authentication-Results: i=2; sourceware.org; dkim=pass (1024-bit key, unprotected) header.d=amd.com header.i=@amd.com header.a=rsa-sha256 header.s=selector1 header.b=2GpRPVgc DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 77B4B4BAE7CD ARC-Seal: i=1; a=rsa-sha256; s=arcselector10001; d=microsoft.com; cv=none; b=gyoTKfi+se80PTrbY2vsebhzTl0QbPCcAybSmDGxdfDeUrDYy6qt5aDaECUJ3Rnjf42jsGPpdM22R75hfBew1hJBzqREV9L87DQngWax4zKAbp2ao+WS2lxBAvgVa0KrnCRPR6dcsgbmnpUMg1UgwpVvgPbys143zar/XxNTn1h+OBeDBxLa8p6bTBzoi4VUrvYB/HHFGrvDO3GCQMVEQ/58rpG4Wh2bbGOdILkoKCiULsfEMKGTIzHP1t4esY497Tj9WFzdqb/SiEFr8NrGyN1dfWitGJ0NnT6mwIwJeHux9+Q+94YztYQk7ejT3mSNHerxoUhUYHplJbdx/7le+w== ARC-Message-Signature: i=1; a=rsa-sha256; c=relaxed/relaxed; d=microsoft.com; s=arcselector10001; h=From:Date:Subject:Message-ID:Content-Type:MIME-Version:X-MS-Exchange-AntiSpam-MessageData-ChunkCount:X-MS-Exchange-AntiSpam-MessageData-0:X-MS-Exchange-AntiSpam-MessageData-1; bh=b/XdMMTInHnQDH555wLC760xx/XpntJj+yUqe1ldPyQ=; b=EUx0aL+urpnPqEVuUvvu0SYT+CFhMdyz7D51JkYmao8k7tWvauFo6t6VX3JlTOdr5SVE5tehqcyb/z081xJe/ZuNqzfVZvO7670LdkEE5tazasK1KtznDUYDmeGDLzEKVY7IFp7rSmiYRKQ5hu/snXDNykbJAypSBrtE5DdnIogRmCjW50mHnbUsi53Mfjfjt41BKIsRf6iHyFqpEXxF5oO1oOujbDeOuQ4AQyiO2fGz+c6tDsAchRYfZj4vneisDwTadJBsclpMC64vtYId6fr3AjSgHqYZXV11HCjRxDl9YqBJlbQMkLcibCpPp9JEi5yaSKzyfFSQ61GEUNwoDg== ARC-Authentication-Results: i=1; mx.microsoft.com 1; spf=pass (sender ip is 165.204.84.17) smtp.rcpttodomain=sourceware.org smtp.mailfrom=amd.com; dmarc=pass (p=quarantine sp=quarantine pct=100) action=none header.from=amd.com; dkim=none (message not signed); arc=none (0) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=amd.com; s=selector1; h=From:Date:Subject:Message-ID:Content-Type:MIME-Version:X-MS-Exchange-SenderADCheck; bh=b/XdMMTInHnQDH555wLC760xx/XpntJj+yUqe1ldPyQ=; b=2GpRPVgcT/NTodnhgU1OH3NypGJrjSjvGPgox+AJ7t73JIjbrimsEWnOX7baFHgzxVq507K8hKl9fDqCv6ULYy+u59woeuD1f8g4coVdJuukoNeCD2VOeV4/QuZhFXA/if8z7121UJZkOWNayC5WFm+gDu9ATslopCtQBi93bRQ= Received: from CH0PR03CA0427.namprd03.prod.outlook.com (2603:10b6:610:10e::30) by CY8PR12MB8065.namprd12.prod.outlook.com (2603:10b6:930:73::16) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_256_GCM_SHA384) id 15.21.292.19; Wed, 5 Aug 2026 22:24:42 +0000 Received: from DS3PEPF0000C37B.namprd04.prod.outlook.com (2603:10b6:610:10e:cafe::81) by CH0PR03CA0427.outlook.office365.com (2603:10b6:610:10e::30) with Microsoft SMTP Server (version=TLS1_3, cipher=TLS_AES_256_GCM_SHA384) id 15.21.292.19 via Frontend Transport; Wed, 5 Aug 2026 22:24:42 +0000 X-MS-Exchange-Authentication-Results: spf=pass (sender IP is 165.204.84.17) smtp.mailfrom=amd.com; dkim=none (message not signed) header.d=none;dmarc=pass action=none header.from=amd.com; Received-SPF: Pass (protection.outlook.com: domain of amd.com designates 165.204.84.17 as permitted sender) receiver=protection.outlook.com; client-ip=165.204.84.17; helo=satlexmb07.amd.com; pr=C Received: from satlexmb07.amd.com (165.204.84.17) by DS3PEPF0000C37B.mail.protection.outlook.com (10.167.23.5) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_256_GCM_SHA384) id 15.21.315.6 via Frontend Transport; Wed, 5 Aug 2026 22:24:41 +0000 Received: from khazad-dum.amd.com (10.180.168.240) by satlexmb07.amd.com (10.181.42.216) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_256_GCM_SHA384) id 15.2.2562.41; Wed, 5 Aug 2026 17:24:40 -0500 From: Lancelot SIX To: CC: , , Lancelot SIX Subject: [PATCH] gdb/amdgpu: Handle SIGABRT with a higher priority than SIGTRAP Date: Wed, 5 Aug 2026 23:24:25 +0100 Message-ID: <20260805222425.588196-1-lancelot.six@amd.com> X-Mailer: git-send-email 2.43.0 MIME-Version: 1.0 Content-Transfer-Encoding: 8bit Content-Type: text/plain X-Originating-IP: [10.180.168.240] X-ClientProxiedBy: satlexmb08.amd.com (10.181.42.217) To satlexmb07.amd.com (10.181.42.216) X-EOPAttributedMessage: 0 X-MS-PublicTrafficType: Email X-MS-TrafficTypeDiagnostic: DS3PEPF0000C37B:EE_|CY8PR12MB8065:EE_ X-MS-Office365-Filtering-Correlation-Id: 2269dee4-76c4-4d76-2cc8-08def3405861 X-MS-Exchange-SenderADCheck: 1 X-MS-Exchange-AntiSpam-Relay: 0 X-Microsoft-Antispam: BCL:0; ARA:13230040|23010399003|82310400026|1800799024|376014|36860700016|13003099007|6133799003|10067099003|11063799006|5023799004|56012099006|3023799007|18002099003; X-Microsoft-Antispam-Message-Info: Yn6M26pBLwYQgMcB/34H6JtkO3Bu502f2lcXQzmlsqKoIHSNBthKBHmtHDsVoMPB7fISSpLgb/ZVRcU09G/3owQ730QXKRMFe5HoxLHlXcie2jEc7q6aBzldpr1I0+CQ53Cx0owpOIawWpSvMuL/hVqju/8dJTX6mjVvpXc1bvmdr5kN/SSUZCfp+da2j3cSoWStp5aZsR0tlL9H52zYCJ5hDPJEXzfB0VMDFrnOkS2SUlczCWJ22hZI9kctiK5QSwj692A4Qt6Jw53IIGWodU9Ha5RqnqnoNFZDu8YomfgzaeOfD+Jga0MXFsBeLdGyfjAbF8KxSLR612HyCGYkGP4ZyOBJP5MkzHMRJ5Z2QkFjVcdl+E6+yBxjLp90s9Lutp0hrfMdUtGOo+b40cqZMoHr2GHgFgS9/P6hALKBBBWj72NLR9LQnn8Kg8apxTxFaCvRbvC1fyZz0W1DepATTQ+w/rNujTep6XtgJDB2/0F6J52jghnBCqAuhxIRREiwyFxEBOmfIF9gZ9KfHf7+zhkNki2Qt+g23eOGw/tE0QzqgzVzV56tDFTHGZc8ghDf9ZE8XhKKvoWKX6tVtP0VXxL4uWvjPz+n/6ASQkvsPAZVZ9WQuGbL08QgRrh9/p8/mwfr02jcxTQlF4eQrhteJmmMDFq+kTeBTF+gPFwvMTtbZ8m0JRDEPLx4b+eaYeizq5nneVGef83R0izvayeUQw== X-Forefront-Antispam-Report: CIP:165.204.84.17; CTRY:US; LANG:en; SCL:1; SRV:; IPV:NLI; SFV:NSPM; H:satlexmb07.amd.com; PTR:InfoDomainNonexistent; CAT:NONE; SFS:(13230040)(23010399003)(82310400026)(1800799024)(376014)(36860700016)(13003099007)(6133799003)(10067099003)(11063799006)(5023799004)(56012099006)(3023799007)(18002099003); DIR:OUT; SFP:1101; X-MS-Exchange-AntiSpam-MessageData-ChunkCount: 1 X-MS-Exchange-AntiSpam-MessageData-0: 7l+Fs9BbyZ8kMTOw2+vWkCzcZ6DqXm5gxRn2KSDeIkWcv+aApBNmSeEnmc+UPQCDsEUp6WykMKAQNWfJuHgl5AEpThJ+qNs0bdr+s2nUh389QNLsTWq2ESBE2LUwjs7TcFukCYppP+RTmaPH1UOdjtalR35HfaoTt79oZNw+PGCsKabfC1Rf7Tda96CvCePdrDs5ymGukorpZJjhabK3TlnRbCwhtCgtDXHmht83GoPzUKsMGAN/6z7dfg2h5oQdAtTRSojzlTCQD7ZMSqAXaYuDXbQYnYPDej28gPQEVh/eM5LL734BlIrT6Vt/3etKsgYQdrNyGRdVFtr/wNN9X6it+xzeXmqph/EdalQz2I6IACAvCW0CbMtnAyEKOer9Lj+AXGovZerno48nix7xQslf59TOiWSmwQFTkAKZPZf+6FW1f304FwnmOOr1PPx0 X-OriginatorOrg: amd.com X-MS-Exchange-CrossTenant-OriginalArrivalTime: 05 Aug 2026 22:24:41.8603 (UTC) X-MS-Exchange-CrossTenant-Network-Message-Id: 2269dee4-76c4-4d76-2cc8-08def3405861 X-MS-Exchange-CrossTenant-Id: 3dd8961f-e488-4e60-8e11-a82d994e183d X-MS-Exchange-CrossTenant-OriginalAttributedTenantConnectingIp: TenantId=3dd8961f-e488-4e60-8e11-a82d994e183d; Ip=[165.204.84.17]; Helo=[satlexmb07.amd.com] X-MS-Exchange-CrossTenant-AuthSource: DS3PEPF0000C37B.namprd04.prod.outlook.com X-MS-Exchange-CrossTenant-AuthAs: Anonymous X-MS-Exchange-CrossTenant-FromEntityHeader: HybridOnPrem X-MS-Exchange-Transport-CrossTenantHeadersStamped: CY8PR12MB8065 X-BeenThere: gdb-patches@sourceware.org X-Mailman-Version: 2.1.30 Precedence: list List-Id: Gdb-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gdb-patches-bounces~public-inbox=simark.ca@sourceware.org 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 . */ + +#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 . + +# 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