From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from simark.ca by simark.ca with LMTP id QLVgJEzvAWrsFDAAWB0awg (envelope-from ) for ; Mon, 11 May 2026 11:01:32 -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=Dp6nKWJT; dkim-atps=neutral Received: by simark.ca (Postfix, from userid 112) id 7E0231E0C3; Mon, 11 May 2026 11:01:32 -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_MSPIKE_H2,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 37B081E067 for ; Mon, 11 May 2026 11:01:31 -0400 (EDT) Received: from vm01.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id A69E24BABF35 for ; Mon, 11 May 2026 15:01:30 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org A69E24BABF35 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=Dp6nKWJT Received: from SJ2PR03CU001.outbound.protection.outlook.com (mail-westusazon11012061.outbound.protection.outlook.com [52.101.43.61]) by sourceware.org (Postfix) with ESMTPS id 17D734BA5439 for ; Mon, 11 May 2026 15:00:46 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 17D734BA5439 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 17D734BA5439 Authentication-Results: sourceware.org; arc=pass smtp.remote-ip=52.101.43.61 ARC-Seal: i=2; a=rsa-sha256; d=sourceware.org; s=key; t=1778511646; cv=pass; b=kVJGBSl48NDlCCNIU7TBprfiWwixUWi0jDH3Wv7358wANqc7F5hIwg9w42TCT+NXRF8G4enb9LY89MRlRQWAn0oVTUPffn0CnxbLfs53zS5+SKpWQsExR2oJZk31hFNOy074HtCurMWpkp84X54mTJMtkZmXVOfNb2zr85GH7Uw= ARC-Message-Signature: i=2; a=rsa-sha256; d=sourceware.org; s=key; t=1778511646; c=relaxed/simple; bh=ZoTjkDnTHBnPI/lXIum9hYHVTYSHeqAFVoWgvsW30Qs=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=McJGilGdMYjpYSFOSng8hQFUX4NwPq8KSYooNvU3QNnVs3N9EiaOII6UX7hFh19GqVEjndnx8ECMlxBLdwjbu/h3niob3ZnKTkq3n4797sVRUdBO5X/C3wX2+INGQtlOkFx4Thm1JQ1nRiHpjlnwvc3GB2obiEFUCZ9P1q310lE= 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=Dp6nKWJT DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 17D734BA5439 ARC-Seal: i=1; a=rsa-sha256; s=arcselector10001; d=microsoft.com; cv=none; b=ueW1nT15C16tCQ1RCHo2NWIxPHEeneo/K/C5KZKtzJymPDaRTYThPFQXsa5fIksQISFezsBFegN2Yk75Xeb+F+tQiFidvSiI0e8GVBKJZITqE7p4vZvYsuByvF0s4V2fY+o4/wz5FITbioXHh9zvA8txbFhZsNjZ2zwOPVtkDfz7gCkzygrj/f+cUGP4dARhFEDrtoXkLlJRnP2k0zMbkK8AI/oMZ5LBheIEM2KZOtwFNAJIbaRNI0P0T3qWMeaJFUvwV3Zs4JNpmF8PCH8vxTtsmNbUuCXxJcrngz6G1afXAsAbJtFWkftyNh4EuYBztOkCylyNrIKuiQyDLJEx+g== 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=8dNPiYdERSVYjBHNX8Uzzbt3leJWknSw9nHHgL+W1cg=; b=e7gyVAS9PoJLpkCOatZC0pXQiHgRIlddC+N4YzENT6LBK2IJroHB060Q4zsBa1S3x5q06+uk6J2wgh+VA8OSLJw2tBdyfvQaizZfkT7GSnl7A04cKwYaspAuqofk2hqWAO/zrwGUwayZEKdKCGpoL41aDN+ZvPQzZLvRHE+OVIr01BVCeQB3kQ0dTQEg17U+XVHEHFcypmtQnxFbnnq11z5yj1lCzxtXFZXxeQrclvahx88hfI24zhmjoNaHW0BVduFwIUH+QKrOjM1lYVhl4mScHvkb7GVTSDHOWZ1n2EPxfHiCBUCmdGrYs0DX9wDeZNDK8bXQNLDvK3o3+ft0sw== 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=8dNPiYdERSVYjBHNX8Uzzbt3leJWknSw9nHHgL+W1cg=; b=Dp6nKWJTdYOOmePvk/28Ihxec2YSMS4kRrwrHBuDXdJSspFi16iX5LpcBjDEdOd3m2KhD1WxPGISFHb36TgUoTqGrrLmScT/pKNuBcRBjVVO0FvU9OVRfv1spa1CbRfoN7d1jZ2qCA9ItICb2/6nxnlHw79jH0H2efUG9ab6tbE= Received: from CYZPR11CA0016.namprd11.prod.outlook.com (2603:10b6:930:8d::21) by PH7PR12MB9126.namprd12.prod.outlook.com (2603:10b6:510:2f0::21) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_256_GCM_SHA384) id 15.20.9891.22; Mon, 11 May 2026 15:00:35 +0000 Received: from DS3PEPF000099E1.namprd04.prod.outlook.com (2603:10b6:930:8d:cafe::43) by CYZPR11CA0016.outlook.office365.com (2603:10b6:930:8d::21) with Microsoft SMTP Server (version=TLS1_3, cipher=TLS_AES_256_GCM_SHA384) id 15.20.9891.23 via Frontend Transport; Mon, 11 May 2026 15:00:31 +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 DS3PEPF000099E1.mail.protection.outlook.com (10.167.17.196) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_256_GCM_SHA384) id 15.21.25.13 via Frontend Transport; Mon, 11 May 2026 15:00: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; Mon, 11 May 2026 10:00:28 -0500 From: Tankut Baris Aktemur To: CC: Subject: [PATCH] gdb/amd-dbgapi-target: implement the 'interrupt' target op Date: Mon, 11 May 2026 09:59:45 -0500 Message-ID: <20260511145945.727267-1-tankutbaris.aktemur@amd.com> X-Mailer: git-send-email 2.34.1 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: DS3PEPF000099E1:EE_|PH7PR12MB9126:EE_ X-MS-Office365-Filtering-Correlation-Id: 26afa617-ef82-495a-4161-08deaf6e0bef X-MS-Exchange-SenderADCheck: 1 X-MS-Exchange-AntiSpam-Relay: 0 X-Microsoft-Antispam: BCL:0; ARA:13230040|82310400026|36860700016|376014|1800799024|13003099007|3023799003|11063799003|56012099003|18002099003; X-Microsoft-Antispam-Message-Info: P/FbjFwrXHU7iEnutKa+/aReFeXD7MrElOLU+FDr39ItzdIdmSJM9Dhpw4O+KaYMUv3AVLuVOSDtQdheFYUtyax/utaSPXNRflZN66N6quOFlLlC+ycESxeGsk6mC1lQtGvTblzim2WnMZ1iLgbzPiIKHV7GjcMFsgxgFdzGGEbeaiqfz44tt7kdZ333trXILPLg8fNBL409MqZrTb/vHL1m2QiJ1/Uki0nLDUeAyepzlaBLX3JKMRtaL3pPqb1lUnEJW3NRz94a4QTOWLUjsYqykSffz26zbNgJvodKkm/7HUKF9ump0mLt4JBt0yn6ZItAco3z4FwI2ZRo+i/RBYBm0Y2aMEAlIt+CiAJHRCO+7hBLQJRKSmdXmq9Oye8laVPl6JgdO7t6+lYfAJwdMvbc/spq+wT889lEMB8fgUPBvNHZWiT0oZ1Dlc4xJdcuGpdo/PmYxZfkkRU85Bqb3DR7lhhYx3Ycgy8tyj2yd6oQXtkTql26pzfCnon24v+Cxx4yEWaaIpkvPV91R7NbBanRPi7vshCypXd8oCnf+LjzMcqQjroUR0On9sJGV3U8gXjjWd0i0Gp5zmMvBqh5hGEmDm2JA1lO89P5J6knwvLo3Cvyj63K1n/V4zw3L6ZIJ+e1Muwd1qMbIfdTgFCZT155gaPXH2CAwnOXnY1Dokzr0LcSeSA9i+xTFlvt7lJj1+d2p/csuw3V44X6TCG9q5Fn3c9sQ/gwHVnXfoQmnH4= 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)(82310400026)(36860700016)(376014)(1800799024)(13003099007)(3023799003)(11063799003)(56012099003)(18002099003); DIR:OUT; SFP:1101; X-MS-Exchange-AntiSpam-MessageData-ChunkCount: 1 X-MS-Exchange-AntiSpam-MessageData-0: 9jcl0daUTNsK0pG/vFI71PU2MaT3twuUz+XUvlYS0ksm6914mMKitd2d1f37ODVSFMUVe+W0cwyrEnXxwffrgnCrfiUmzcxI7atdqtGintjo1aIh0j2AwqwSiFwWbRKHxyItdK0XQQp92LHyz6hnPbTXSnBcCL75IJjpooAJXoWmphq+qk8fUXJeY5x7JXTQzWpd1ObJWaWAk2n8HeGsldDfF4hwLZEfhCOu374Y+hOhQFRNzv6pwOujvIgngMqpL+JiK7LsY4LLuoO9wcgDiMqELvKFaubeQ6uoEecqsVhnM/KYvHvrSGl4eK80TFoxGaq/pjjROC8QNL5IlUgP14L9TTZXhS+U71rWja4oV4OJTcIfaIw3P0Jd0M4bc3FXhSdROqYes2RZVDgTMSQmif5hO5vBOX7/g2Q4yLvPicBBv3j0qAmyCrIRIV+ha6yr X-OriginatorOrg: amd.com X-MS-Exchange-CrossTenant-OriginalArrivalTime: 11 May 2026 15:00:31.4072 (UTC) X-MS-Exchange-CrossTenant-Network-Message-Id: 26afa617-ef82-495a-4161-08deaf6e0bef 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: DS3PEPF000099E1.namprd04.prod.outlook.com X-MS-Exchange-CrossTenant-AuthAs: Anonymous X-MS-Exchange-CrossTenant-FromEntityHeader: HybridOnPrem X-MS-Exchange-Transport-CrossTenantHeadersStamped: PH7PR12MB9126 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 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 (); +} + +int +main () +{ + /* Make sure that if anything goes wrong, the program eventually + gets killed. */ + gdb_watchdog (30); + + kern<<<1, 128>>> (); + 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 + } + } +} -- 2.34.1