From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from simark.ca by simark.ca with LMTP id 6YOEHgwoA2pxjzMAWB0awg (envelope-from ) for ; Tue, 12 May 2026 09:15:56 -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=pfkGVqCp; dkim-atps=neutral Received: by simark.ca (Postfix, from userid 112) id 70FCE1E0C3; Tue, 12 May 2026 09:15:56 -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 65E611E093 for ; Tue, 12 May 2026 09:15:54 -0400 (EDT) Received: from vm01.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 338D54BAD175 for ; Tue, 12 May 2026 13:15:53 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 338D54BAD175 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=pfkGVqCp Received: from SA9PR02CU001.outbound.protection.outlook.com (mail-southcentralusazon11013028.outbound.protection.outlook.com [40.93.196.28]) by sourceware.org (Postfix) with ESMTPS id E54394BB1C01 for ; Tue, 12 May 2026 13:15:09 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org E54394BB1C01 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 E54394BB1C01 Authentication-Results: sourceware.org; arc=pass smtp.remote-ip=40.93.196.28 ARC-Seal: i=2; a=rsa-sha256; d=sourceware.org; s=key; t=1778591710; cv=pass; b=g2mnacNofnWLu69OqYnauSG9DFQSR2kgGSrX4mkosQqbUviWXVnreRzEsAynNtjpMfucXI6x+r2wVYmDqmGJ/ywtJsov53qm93cPvA93D6+hNtsF61nJwjzazxVQb6f4g5h0LfTgM6ApjqvDekHHO/ymUv3ECCV8fgaU2AkB/98= ARC-Message-Signature: i=2; a=rsa-sha256; d=sourceware.org; s=key; t=1778591710; c=relaxed/simple; bh=sxfGlkHAI2MI2MCp0+TbhT1cejEIkIdWtj/8p8lzz2Q=; h=DKIM-Signature:Message-ID:Date:Subject:To:From:MIME-Version; b=nfWI/9q53eHJD7DkLSQ6CVEcOchqZAf93eq/bTN0BqWLoEtKUnCjoUSsOgh0wrnrLE3KfFfxAoJNtmAMKIHZEfhNGFPaHx798MXyYK+BKwkjfXNgX2CQ6GIzqX+m6Z8LAazZiu5UpfeKDwAFpBR26txYSHbpKM3iFw9TBgrRkn4= 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=pfkGVqCp DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org E54394BB1C01 ARC-Seal: i=1; a=rsa-sha256; s=arcselector10001; d=microsoft.com; cv=none; b=EpJn5l/hR8p/KREzoDZ1eOEhGJQqKikRPt4crBj71P7cdoCBlmcspB+GjDAC5cwY+rWHRu7UE11x8MuZgmTtzaZLI7Tq3l777EVF0Q0x79TPDlC7XAJs52LtKcXKjHRelwJULmJbeijIl8Yz0R3X0TblzQ1M8BN8A4OyeA1haUzpFrCRW+t0aisVzCjoTrwm8kIDzLg5C1BvcSmOXtILZdC6aWQi87VtwKa7OBcGZQAcqFwmk8MVBA5M45FFSpOjs4UyNdSS4g/EwnMSAL5S/OqVGYc7BfxkEvJWDn1aqARKxj6KBa1tAxL8+hwpuPKZyQ1crLTfoG7rfLUC9O992A== 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=DgNoBjcKSnj4qkkD1a+cRv3pjlWLzP8gZY3ZsKM9U9c=; b=WlPeSylLhQRnafE7fSv+aNXB4tu5mzdBai3MWb3vh1nM5En3Ww6ZWOA7BuAMRgtjzr8EYsGO4Jn7cWq2DyKiUl8QP2fnL7b9R9s9EQvrJR1RIeUfBAHtb5qL/+8XSZHH1e7H7EFRMk9d5LcO9WQEIyWk2v5z/mDAEpUp4EJuFyWdiyKLN+B0da5uRdIJNyu2llTNAE32pbP1MSqSNiSHSE4dQ/pAVVN+yxUZ3ZdM0NQDWPkVmKSiKJAZHu5POYHokWIjKjqymCBdmZEbCbvvPuvEqR2JUlJyxelmrF8yDHspd0mcxhOqpC9B4BsMQigahSSqaDDiyTRGOmxJeexjnA== ARC-Authentication-Results: i=1; mx.microsoft.com 1; spf=pass smtp.mailfrom=amd.com; dmarc=pass action=none header.from=amd.com; dkim=pass header.d=amd.com; arc=none 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=DgNoBjcKSnj4qkkD1a+cRv3pjlWLzP8gZY3ZsKM9U9c=; b=pfkGVqCpuv5UDPjHLKTTrEfwbxXCn+vNl4Nd3lIvywPhBAjWQomW278Q0OG7L+Yveze00ckGisuNwWL0GhRYhZnB+aKu+G4agBvR6XhVXdRtzVRRD+Xadza9K8rYgyT4O0ek15r5ckUg/NOxYvLUIkf304Cyrp7jThh/RI18Mi8= Authentication-Results: dkim=none (message not signed) header.d=none;dmarc=none action=none header.from=amd.com; Received: from CH3PR12MB9079.namprd12.prod.outlook.com (2603:10b6:610:1a1::9) by EAYPR12MB999132.namprd12.prod.outlook.com (2603:10b6:303:2c2::11) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_256_GCM_SHA384) id 15.20.9891.23; Tue, 12 May 2026 13:15:05 +0000 Received: from CH3PR12MB9079.namprd12.prod.outlook.com ([fe80::5a7a:2adf:1862:35c3]) by CH3PR12MB9079.namprd12.prod.outlook.com ([fe80::5a7a:2adf:1862:35c3%5]) with mapi id 15.20.9913.009; Tue, 12 May 2026 13:15:04 +0000 Message-ID: Date: Tue, 12 May 2026 14:14:59 +0100 User-Agent: Mozilla Thunderbird Subject: Re: [PATCH] gdb/amd-dbgapi-target: implement the 'interrupt' target op To: Tankut Baris Aktemur , gdb-patches@sourceware.org References: <20260511145945.727267-1-tankutbaris.aktemur@amd.com> Content-Language: en-US From: Lancelot SIX In-Reply-To: <20260511145945.727267-1-tankutbaris.aktemur@amd.com> Content-Type: text/plain; charset=UTF-8; format=flowed Content-Transfer-Encoding: 7bit X-ClientProxiedBy: FR2P281CA0135.DEUP281.PROD.OUTLOOK.COM (2603:10a6:d10:9e::19) To CH3PR12MB9079.namprd12.prod.outlook.com (2603:10b6:610:1a1::9) MIME-Version: 1.0 X-MS-PublicTrafficType: Email X-MS-TrafficTypeDiagnostic: CH3PR12MB9079:EE_|EAYPR12MB999132:EE_ X-MS-Office365-Filtering-Correlation-Id: 3a0a45a2-66e0-4cec-8f2e-08deb0287b1f X-MS-Exchange-SenderADCheck: 1 X-MS-Exchange-AntiSpam-Relay: 0 X-Microsoft-Antispam: BCL:0; ARA:13230040|1800799024|366016|376014|13003099007|22082099003|18002099003|3023799003|11063799003|56012099003; X-Microsoft-Antispam-Message-Info: CAuGOHzmYtLFdBg5ZMI98WkXaJDsFcWdxvcQfZX7sssXfaDTBwLycMOb3UQEKqoIEG6g1M0WCwB7XcYTE0YpZOs3dXKxowxISfvST5Zz7x6XtI0Z3nqMVmtpDxKafgTf+oQcgQkxRHF3JH8yVXiRWrR8osirhxbgjYW3Glap3Uo9YqOIvXMhdxUBD1h1ycr2VPxleOevcdhxYI3laOw151Mz9lZCGOR/e+t8N6VkEiJUmUH7JxSRao6lcaSeazFxTImVTvvT76a41VPqJ+Yi3lYMnuAdSE5MvRbA+nyid08IUIbECopcBp1pt6yqvUxlX9wjrrYCzmGHW07m2B9Bh3kgTqxKKYsXEjz32XCaqqwTS8VbEKVwZOtk8NeAN1IUA6Q4ir2C4EaTNpd9Kk1CQqlezzUWd1GpcfabAYrDi2ZJGtEucdwXsKCubkpn0YUg6Qvrg47Ctuv+S4VraTPfGd42bRlTVipCXhWbWqjkXNTQDtAfgmr+R/6Pt1udQD/DLtRuraBLLw4uZEDunUeYq3xb0Vdb2a9tFGmsSjGnNEim6P1bCnEuzsuum+TS6+bbkwYMRTKtEPaSOhQrr4tnKomIg0qL49dC945YiCcPvO3zZ5s32AuJCzRfDPvbdiMbygLM1xyWIJTEj24yT4LvNLLp+Jmfhw88WA9n1RwkKrzlUfEIuq3RQqFKmF91On4JPiiRLF6O/GJllZZ8mdtrDA== X-Forefront-Antispam-Report: CIP:255.255.255.255; CTRY:; LANG:en; SCL:1; SRV:; IPV:NLI; SFV:NSPM; H:CH3PR12MB9079.namprd12.prod.outlook.com; PTR:; CAT:NONE; SFS:(13230040)(1800799024)(366016)(376014)(13003099007)(22082099003)(18002099003)(3023799003)(11063799003)(56012099003); DIR:OUT; SFP:1101; X-MS-Exchange-AntiSpam-MessageData-ChunkCount: 1 X-MS-Exchange-AntiSpam-MessageData-0: =?utf-8?B?VkRSQzRsejBGUWVKTi9RR212NjduRTZvMnBnU1ZWTXFSWGZFaUQwYUdJWjFl?= =?utf-8?B?dDhnTXY0ZUpzTUxkbk5EdlB0WUcraXVyVkhCU0V6L29CTzhyaWI3Rk9pYk53?= =?utf-8?B?MmhrclhwdDZUeURhbVF0MnhnaUNBZnYvZXEwbFgyMWdUaG5ZMmFDZnROUjJk?= =?utf-8?B?bnRwR3hDbzlGQmVKbDlDTGJKUkFwbnNqb0xiTTZvbDkxWU5jQ2pIUUlVSHBy?= =?utf-8?B?bjZtc05WdG9laHlrYlZQaCtIcE5tZFJEd0FITXdrMmRrZU50dDFmYW9ramZX?= =?utf-8?B?Ly92UEFrV0V0NEdDTGxmTDlqMmVOejI0VlJVSHNuS3J5TWtwQ2ZNSFNmQ0po?= =?utf-8?B?THZwS0JjZ0VmTzREM1dBYTAwSWxxd2UzZmhQRTlqallFK2d4b1V0NVpJcUFn?= =?utf-8?B?SytRYnV3U0o1dFM5bGRKM3JRaGtqYlNEUWhQenZpYVVkMjNzNkozRWpqV3R5?= =?utf-8?B?am1PRTZ1UHJUR1piWEZHZnFHdmRud0ZEOGZZZGQ5YVgwbmRMbzJ0OG10ZStR?= =?utf-8?B?ekU4a2Nxa2FkdGNaeThBakJmQWw3N1YyUGJSQXR6SHRWWXBpZmJaeG1GVGRI?= =?utf-8?B?Vk9yRC9sSCtzZHlGN0lyUVczTUdtRy9tMzNkbEhQemw3UjdCTkJ4T1prSDAy?= =?utf-8?B?T2xlQkljK0RrMmltbFBxbjVCMWRyNkhXMmFwS1ZaWStqM1dWcWhuUVlwOXZC?= =?utf-8?B?U3BOM1kxWkhhWjlKMjg0MTNkOFRFOWJWSUtPdldmaDJzTjdJQ3BGQ1NwRWtY?= =?utf-8?B?UXN2RU44VkgzV2k3Q2dUK1E1SUw1RktZMFd1aUoyMHhWNWllemdETkU1VnFl?= =?utf-8?B?Z1ZibSthakpTdStEVGhGMXMzdWpMN0t6VXNoeXY2WGpWV0todHZDUHBNazcr?= =?utf-8?B?T2trTVIvd1ZvdGFZR3hMMDFqM1FNSnU3MFI1ZHdSRFR1QlZqRTlFdlhxT21z?= =?utf-8?B?RldXZ09hODVmOFNueDJuM0p6cHdFN0dYamxEMXY5b1I5UGV2YjlpaTdwZytG?= =?utf-8?B?M2RvbktIdUptVlN6cTErb0M1U0xpTis4MGRic2JOaW4zcGhzM0M3YkErZkQ3?= =?utf-8?B?YjI1NHdWdGpsUktlbVB6amlRWTdyUDdFWXVHWkJrZm44UHhTTGE2Y2Nnem5X?= =?utf-8?B?VFBNTjkxUm8wM3hHYTNldG1iRkJFaWJpNmJUd2lQVzMzZVc3VWxDd1pHV0dL?= =?utf-8?B?ZngwNnlxMkc3dGgxc2pPVHhxN3RUN2p4ZlV4dy9QZDNUOUhDOC9qa1ZzaEtr?= =?utf-8?B?MnV3SlhLRU5OYUpFMEU5SUpJWHZMM0RCRC8vR3V4c3NlUkp5dDM4RzIxcVFz?= =?utf-8?B?T1lVVVUrblVJSjNnSDkrQmpQTlBicGRTRnBPOGJuclJpaGJXT0lzV3AzUzhv?= =?utf-8?B?MnNsL2lwZER0bXFMQ05lR0s1WEdXblpLZGhBeUdNZGNCVmhTVjJoZnJERXY2?= =?utf-8?B?YzFyaHpjMXI1MnFCb2U3Y2R2ZlhHZHhLNUg5eFRNT0lPd3l5dDZWQ0Ntam96?= =?utf-8?B?MlZmOHg1bjlIWjR4U3NnZWRmd0x3bjhTcmVoaXRmNnF3YTI1bGhMLzRsTlRJ?= =?utf-8?B?b3NBV0ZxTDMyS0pPZVpxR3p1QzJyekFFeEJpZkJ6N0NlV3NJSUx3MUVLUUxj?= =?utf-8?B?ZXZJQ3NzbnRHclpWOUJVSmV3QmJDWG54WVlQelJQNC9WTlMvTDBNaGVLcURD?= =?utf-8?B?U01MaDkvVlgrVU5nQ1ZQZEFGa0VHS1pJdkxTK2tkd3hGVGd6MTZJTzhjdWZM?= =?utf-8?B?eWlreHBNekJOUG9FNldMN0NEV0RFaTM2bnlOREdZUzlZSndQZmFpZ3A0NzJH?= =?utf-8?B?MGhYTkNpRXV5bk9iU1VpVTdsNW8vciszUXRNeGRIYmRQTmZwb3F2QnpRNXNi?= =?utf-8?B?UDJISGhlV0tJdzIyR0NFUklnOE13ZFN3NUM4S1pkaS8yUkM4aUIxVHBiUlZ2?= =?utf-8?B?blMxUXo2WitkUVBrcjJTYVdjRFVFSllQbi9PaDh6OVA3VEJNMVZqcDJ6aUJJ?= =?utf-8?B?eitmTGhjYkJNRlJ5TG1xZDhlN1VSdm9pcVNGVlpGVEMyQ09paXdEd1Fic1g2?= =?utf-8?B?dTg3NkhaRHEzR28zQnJkdnFjc09qeEZCeTlRTmVoZkVIWnpYZHZraGNweXBC?= =?utf-8?B?S1AxWGNLcWE4VWt3QWtKRWwwRjQ3N0J2RDhaR1RGUGVqK2VWNVlOTm9CVEVz?= =?utf-8?B?YlJ1RmRhN2V3QU1nNnRValRoVkI5aDV6YVdzbFFJVmpJbXBudTZ2U3JEejVN?= =?utf-8?B?VUEvRTZsaWVncW83bU5NcTRVQWIyMkY1OGx0RUozUE00WTV1WnlrWmgyTjBK?= =?utf-8?Q?tfLq6V/0BDA9bNw2sf?= X-OriginatorOrg: amd.com X-MS-Exchange-CrossTenant-Network-Message-Id: 3a0a45a2-66e0-4cec-8f2e-08deb0287b1f X-MS-Exchange-CrossTenant-AuthSource: CH3PR12MB9079.namprd12.prod.outlook.com X-MS-Exchange-CrossTenant-AuthAs: Internal X-MS-Exchange-CrossTenant-OriginalArrivalTime: 12 May 2026 13:15:04.6797 (UTC) X-MS-Exchange-CrossTenant-FromEntityHeader: Hosted X-MS-Exchange-CrossTenant-Id: 3dd8961f-e488-4e60-8e11-a82d994e183d X-MS-Exchange-CrossTenant-MailboxType: HOSTED X-MS-Exchange-CrossTenant-UserPrincipalName: Dig2IEopT6ikkfNXCkMKkGMT5WIL1rXJ7l8rKbNJAyQkt8Jt6bO4q+7oebS1wTScNLJfEGs3IHU7s0GZhv1gEA== X-MS-Exchange-Transport-CrossTenantHeadersStamped: EAYPR12MB999132 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 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 . */ > + > +#include > +#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. > +} > + > +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. Otherwise, this looks good to me. Thanks. Approved-by: Lancelot Six (amdgpu) > + return hipDeviceSynchronize () != hipSuccess; > +} > diff --git a/gdb/testsuite/gdb.rocm/interrupt-single.exp b/gdb/testsuite/gdb.rocm/interrupt-single.exp > new file mode 100644 > index 00000000000..fbe357abf1b > --- /dev/null > +++ b/gdb/testsuite/gdb.rocm/interrupt-single.exp > @@ -0,0 +1,53 @@ > +# 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 a running thread can be stopped with the 'interrupt' > +# command. This is done in all-stop mode to ensure the 'interrupt' > +# target op is used. > + > +load_lib rocm.exp > + > +require allow_hipcc_tests > + > +standard_testfile .cpp > + > +if {[prepare_for_testing "failed to prepare" $testfile $srcfile {debug hip}]} { > + return > +} > + > +with_rocm_gpu_lock { > + if {![runto_main]} { > + return > + } > + > + gdb_test "with breakpoint pending on -- tbreak loop" \ > + "breakpoint $::decimal \\(loop\\) pending." > + > + gdb_test "continue" "breakpoint \[^\r\n\]+ loop.*" > + > + # Resume a single GPU thread asynchronously to be able to use > + # the interrupt command. > + gdb_test_no_output "set scheduler-locking on" > + gdb_test "continue &" "Continuing\." "continue async" > + > + gdb_test_multiple "interrupt" "" { > + -re "Thread $::decimal stopped" { > + pass $gdb_test_name > + } > + -re "$::gdb_prompt" { > + exp_continue > + } > + } > +}