From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from simark.ca by simark.ca with LMTP id pqCHDqKgDWowMwoAWB0awg (envelope-from ) for ; Wed, 20 May 2026 07:53:06 -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=FrwfgxrP; dkim-atps=neutral Received: by simark.ca (Postfix, from userid 112) id 284181E062; Wed, 20 May 2026 07:53:06 -0400 (EDT) X-Spam-Checker-Version: SpamAssassin 4.0.1 (2024-03-25) on simark.ca X-Spam-Level: X-Spam-Status: No, score=-3.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,RCVD_IN_VALIDITY_CERTIFIED_BLOCKED, RCVD_IN_VALIDITY_RPBL_BLOCKED,RCVD_IN_VALIDITY_SAFE_BLOCKED autolearn=ham autolearn_force=no version=4.0.1 Received: from vm01.sourceware.org (vm01.sourceware.org [38.145.34.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 51C691E062 for ; Wed, 20 May 2026 07:53:05 -0400 (EDT) Received: from vm01.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id D77304BB5903 for ; Wed, 20 May 2026 11:53:04 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org D77304BB5903 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=FrwfgxrP Received: from PH8PR06CU001.outbound.protection.outlook.com (mail-westus3azlp170120001.outbound.protection.outlook.com [IPv6:2a01:111:f403:c107::1]) by sourceware.org (Postfix) with ESMTPS id DF4CC4BB3B91 for ; Wed, 20 May 2026 11:52:38 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org DF4CC4BB3B91 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 DF4CC4BB3B91 Authentication-Results: sourceware.org; arc=fail smtp.remote-ip=2a01:111:f403:c107::1 ARC-Seal: i=2; a=rsa-sha256; d=sourceware.org; s=key; t=1779277959; cv=fail; b=vqVHEClJCdjdcSF/jsB29o7ppoN3Ai185H8ak5D3q3mqWlVTB4g3L7uErHidjqOKlhSjeiPpuO9tZd1lVL6bSkR+5A2egnfHo7304NbaSGC+Ew02Joz8FeoS79B/BsmoMliqMeVz29S5dKQPN19eLpp04Ht2jxhwVqQwbLZFhPY= ARC-Message-Signature: i=2; a=rsa-sha256; d=sourceware.org; s=key; t=1779277959; c=relaxed/simple; bh=4ReDe30Yg224pFqnKjjSF29lJ4qDz5jtdzAK9voPpOw=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=sfwdDXFmaTDghxfFp9ao5pF2ERdO+G101XMS9cNB6t1FAmX/9+68Pqw+79B+5BV7AVz4dKH5cH8Z3ekoQeyTF0foARWUGHDufU0T4hXcgD+sXY4jsoOwSMx9CjPtSx9b5SSo0jUV4vWCq8Oft5QtB/TpDmmO01n4QSX8m/wot18= 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=FrwfgxrP DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org DF4CC4BB3B91 ARC-Seal: i=1; a=rsa-sha256; s=arcselector10001; d=microsoft.com; cv=none; b=E4jpTd40Xltz91U+3irl4yn2aV/dQ7klg94StGL4YEEfw1/ffeF49LbJ/usfXc46WTW88Ahus26XOXozT+2f8/6mU4ykLhQg7+UqPHapmCIPvXeyP1PL2Ze/HTnsypmR7OyRe2rFEgKhu4K915cxaJKC/u+4Z+7M7DFAuqSX+5CfN5Buj4jkF6WLD5GvXNBNcGn0yDygpe1aJJw/EKVY1QIcN209VanB1YUJY7vElBGqvL0EPgtCTiBopXK2DHg40IS53npNXQVO6PRfZROONvep6lJuF92Ww2K+PLHMxtLxJhVa0YgOZzE20U/InZ0C/FI93/npMeqDtNh6f22NGA== 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=0g/NIJ4YK+ezJXQyUfma02EwsXg4xT6nsyTijqgnvug=; b=le+Kk1frg+CkGrvEHjQ9TDIYazGNlCsQVzJAOtScxWfO2DptukP+XAe5ZKow/4dThPmyFQrQiBxxOm2hALFky1DjAXIh/k5qx/5tuXAwAKlu2lm/7IbOj3za6MF/Ux/1vqoYg0mSJ8JHMRd/zTAyrkyB8J4BNIHq66FmzkjjpU1ca646KpPwUNjobIoLU7igWar8yBAdtPpCHRGP9s3vVhpYZhFcGlmXME+0NJo/tGploOcKnRcxmZbTPRtVLbxdB8VTO8EAVvwKthe+al6z9SUsj++wE6Pw6OOwlCi8VnHkfvf1A0H8m16I9ujOzAgyGLohIp46d6prVPeSnLu6Ww== 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=0g/NIJ4YK+ezJXQyUfma02EwsXg4xT6nsyTijqgnvug=; b=FrwfgxrPy6p/rhaPLvqBz67OoiUXLvvIdRSCD9tQOMgSWXTy/8pbdnDi1RciGmeKjxw9QMENj1i5KUJ418uA6DAi/95TlUQ8VyUEL7LFWScb3A/LnJqoBxGP8G05xM4TRRdMjm/BJOrYNgKkcQaEMG5VHa1ZdFItlwgYDM95Iy4= Received: from BL0PR02CA0006.namprd02.prod.outlook.com (2603:10b6:207:3c::19) by DM6PR12MB4298.namprd12.prod.outlook.com (2603:10b6:5:21e::9) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_256_GCM_SHA384) id 15.21.48.14; Wed, 20 May 2026 11:52:32 +0000 Received: from BL6PEPF0001AB4F.namprd04.prod.outlook.com (2603:10b6:207:3c:cafe::ce) by BL0PR02CA0006.outlook.office365.com (2603:10b6:207:3c::19) with Microsoft SMTP Server (version=TLS1_3, cipher=TLS_AES_256_GCM_SHA384) id 15.21.48.16 via Frontend Transport; Wed, 20 May 2026 11:52:32 +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 BL6PEPF0001AB4F.mail.protection.outlook.com (10.167.242.73) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_256_GCM_SHA384) id 15.21.48.11 via Frontend Transport; Wed, 20 May 2026 11:52:31 +0000 Received: from ctr-rack32-mi300x-1.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, 20 May 2026 06:52:31 -0500 From: Tankut Baris Aktemur To: CC: Subject: [PATCH v2] gdb/amd-dbgapi-target: suppress a repeated stop request Date: Wed, 20 May 2026 06:52:15 -0500 Message-ID: <20260520115215.2042499-1-tankutbaris.aktemur@amd.com> X-Mailer: git-send-email 2.34.1 In-Reply-To: <20260520085820.1299345-1-tankutbaris.aktemur@amd.com> References: MIME-Version: 1.0 Content-Transfer-Encoding: 8bit Content-Type: text/plain X-Originating-IP: [10.180.168.240] X-ClientProxiedBy: satlexmb07.amd.com (10.181.42.216) To satlexmb07.amd.com (10.181.42.216) X-EOPAttributedMessage: 0 X-MS-PublicTrafficType: Email X-MS-TrafficTypeDiagnostic: BL6PEPF0001AB4F:EE_|DM6PR12MB4298:EE_ X-MS-Office365-Filtering-Correlation-Id: 2157511e-2524-4901-f254-08deb6664659 X-MS-Exchange-SenderADCheck: 1 X-MS-Exchange-AntiSpam-Relay: 0 X-Microsoft-Antispam: BCL:0; ARA:13230040|36860700016|376014|82310400026|1800799024|13003099007|56012099003|22082099003|18002099003|5023799004|11063799006|3023799007; X-Microsoft-Antispam-Message-Info: XW8BnNSpzbpZyk8uj2++P250U3LK4LRWPQgEG2TrrdNkpl7XzMCgszp6sRdc9e2EXpQFUC8ieh8XOxRvf6x7LemeQYMe9xCKB8v7iwQjmcnjZt9oFkEoFjCURZtqbXgF6ljxQ4ZdUVzYQJfeO3gNvmPeyBRSCOU37MHiYAe/aPwruJokxtyDqEcTrDIjHRf6MEV54QXgi1OLEkV1FSQCuRVSSv/pxESmAH6MCyw5VSJu0cjMGaowpvTV9heP5h6ljdRaSb9KHj5IKTUEvv8LocOuywmcgNQPur/TVHQuxIKpgkgI/mrzMBGyIdHqX5HCHksbvV0x4ACas3g3zaVq74GtisrXBfcf0aDxqskd2QO8X9WCeokUq0MdNOZf5eDlWxxK3L8x+O5eSnbYOF4Lh2i0leslfsZJMgGPUg9IevkIwtPGZFXFDIKNIsWCQnR7I4ZlsLvbugr2VDtuyy2+MFSZyiONeU8A/yeH5upmxa5xNJ8jd8WkLc74cXffUQ6A2DwRO7JMZyCW74PkVkX7isWZ3Vi7qmZmYxg3ckEgN1yKhMNAMeW4iAF2NCNLOFlSllfXfA2/X0AA6uHTqhjul6OMrNgC2yixE+yNmlQxDU2qn3CFRQJ1ObBClrSwnw90i8cGambwWDseHvOJz2CbA0RawdFZ590RKFY2arVt/74n6c5FbZX5ltD80Fm1K8dBrM9mDPZ8G5ZTXv14rZlHnb3ZF9qXBb9q5KfJPCn68GE= 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)(36860700016)(376014)(82310400026)(1800799024)(13003099007)(56012099003)(22082099003)(18002099003)(5023799004)(11063799006)(3023799007); DIR:OUT; SFP:1101; X-MS-Exchange-AntiSpam-MessageData-ChunkCount: 1 X-MS-Exchange-AntiSpam-MessageData-0: BSzkiv4WFTrFrHWlelsxu6o/+Lo39YkBcxU1ZjGJ9NIZHvNBXRVatE/u0czJ3yzlhuGQD/hWzJpZI4Pw80QJlwt9Vmum2Wqp4hszzBDqOKuq4OGLk404m2P0hF/38CJMvnimRFBMp3gN9jhDTA2Fl2nnGmmmsqAProhEXVNrOR9zDP1oKRM2EEtdKjupmrBjJeWJiS48bS/0ZCQox9ICsp+mfaa/+5vD0rW9RHLb257UNmdGi45GKLRwYvK3+++N8xwba/uCDac3IFkAUHDRlDMNhMwLQt7ofuiH4gBogthKhikZGzhxXEKIhr25R4PeM7CfO5S70YayKsbHELf4Qyu3bCpmTgs+ZVe7s4IPCBw7YCqXnNGWiBiDgI2pBogGETEkJJJxTzIHu/IsNPzUOcjK0s8Dyp1hPSGHIlZkWDwexKfswXi8kWqhC5vepbZS X-OriginatorOrg: amd.com X-MS-Exchange-CrossTenant-OriginalArrivalTime: 20 May 2026 11:52:31.6058 (UTC) X-MS-Exchange-CrossTenant-Network-Message-Id: 2157511e-2524-4901-f254-08deb6664659 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: BL6PEPF0001AB4F.namprd04.prod.outlook.com X-MS-Exchange-CrossTenant-AuthAs: Anonymous X-MS-Exchange-CrossTenant-FromEntityHeader: HybridOnPrem X-MS-Exchange-Transport-CrossTenantHeadersStamped: DM6PR12MB4298 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 This revision uses a nested gdb_test_multiple/gdb_test when defining the user command in the test. Regards, Baris ==== Sending a second stop request to an AMD GPU thread before fetching the event caused by the first request leads to an error: wave_stop for wave_1 failed (The wave has an outstanding stop request) Prevent sending a new stop request if there already is an outstanding one. The fix is in amd_dbgapi_target::stop. A regression test is included. The test uses non-stop mode and executes the "interrupt" command twice, because in non-stop mode this command uses the 'stop' target op, where the fix is applied. To be able to execute two interrupt commands repeatedly, we define a user command. --- gdb/amd-dbgapi-target.c | 8 ++- gdb/testsuite/gdb.rocm/interrupt-twice.cpp | 43 +++++++++++++ gdb/testsuite/gdb.rocm/interrupt-twice.exp | 75 ++++++++++++++++++++++ 3 files changed, 123 insertions(+), 3 deletions(-) create mode 100644 gdb/testsuite/gdb.rocm/interrupt-twice.cpp create mode 100644 gdb/testsuite/gdb.rocm/interrupt-twice.exp diff --git a/gdb/amd-dbgapi-target.c b/gdb/amd-dbgapi-target.c index 421ec8599ed..d44f03d0b80 100644 --- a/gdb/amd-dbgapi-target.c +++ b/gdb/amd-dbgapi-target.c @@ -1090,14 +1090,16 @@ amd_dbgapi_target::stop (ptid_t ptid) sizeof (state), &state); if (status == AMD_DBGAPI_STATUS_SUCCESS) { - /* If the wave is already known to be stopped then do nothing. */ - if (state == AMD_DBGAPI_WAVE_STATE_STOP) + wave_info &wi = get_thread_wave_info (thread); + + /* If the wave is already known to be stopped or there is an + outstanding stop request, then do nothing. */ + if (state == AMD_DBGAPI_WAVE_STATE_STOP || wi.stopping) return; status = amd_dbgapi_wave_stop (wave_id); if (status == AMD_DBGAPI_STATUS_SUCCESS) { - wave_info &wi = get_thread_wave_info (thread); wi.stopping = true; return; } diff --git a/gdb/testsuite/gdb.rocm/interrupt-twice.cpp b/gdb/testsuite/gdb.rocm/interrupt-twice.cpp new file mode 100644 index 00000000000..fc8d2cca697 --- /dev/null +++ b/gdb/testsuite/gdb.rocm/interrupt-twice.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 . */ + +#include +#include "gdb_watchdog.h" + +__device__ void +loop () +{ + while (true) + __builtin_amdgcn_s_sleep (8); +} + +__global__ void +kern () +{ + loop (); +} + +int +main () +{ + /* Make sure that if anything goes wrong, the program eventually + gets killed. */ + gdb_watchdog (30); + + kern<<<1, 1>>> (); + return hipDeviceSynchronize () != hipSuccess; +} diff --git a/gdb/testsuite/gdb.rocm/interrupt-twice.exp b/gdb/testsuite/gdb.rocm/interrupt-twice.exp new file mode 100644 index 00000000000..3c653547dc2 --- /dev/null +++ b/gdb/testsuite/gdb.rocm/interrupt-twice.exp @@ -0,0 +1,75 @@ +# 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 . + +# Test that sending repeated stop requests to a running GPU thread +# does not cause a failure. This is done in non-stop mode because +# "interrupt" command in this mode uses the 'stop' target op. + +load_lib rocm.exp + +require allow_hipcc_tests + +standard_testfile .cpp + +if {[build_executable "failed to prepare" $testfile $srcfile {debug hip}]} { + return +} + +with_rocm_gpu_lock { + save_vars { ::GDBFLAGS } { + append ::GDBFLAGS " -ex \"set non-stop on\"" + clean_restart $::testfile + } + + gdb_breakpoint "loop" {allow-pending} {temporary} + gdb_run_cmd + + set gpu_thread "undefined" + + gdb_test_multiple "" "hit breakpoint" { + -re -wrap "Thread ($decimal) \[^\r\n\]*hit Temporary breakpoint.*" { + set gpu_thread $expect_out(1,string) + pass $gdb_test_name + } + } + + gdb_test "thread $gpu_thread" "Switching to.*" "switch to gpu thread" + + # Resume the thread in the background. It will loop. Then we + # interrupt twice. To be able to run the "interrupt" command back + # to back, we define a user command. + gdb_test "continue &" "Continuing." "continue async" + + gdb_test_multiple "define inttwice" "" { + -re "End with .*\r\n>$" { + gdb_test "interrupt\ninterrupt\nend" "" $gdb_test_name + } + } + + # For logging purposes. + gdb_test "show user inttwice" + + gdb_test_multiple "inttwice" "interrupt twice" { + -re "wave_stop \[^\r\n\]+ failed \[^\r\n\]+ outstanding stop request\\)\r\n" { + fail $gdb_test_name + } + -re "Thread $gpu_thread \[^\r\n\]*stopped" { + pass $gdb_test_name + } + -re "$gdb_prompt" { + exp_continue + } + } +} -- 2.34.1