From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from simark.ca by simark.ca with LMTP id sAoaDL1kA2qj7jMAWB0awg (envelope-from ) for ; Tue, 12 May 2026 13:34:53 -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=h7c/F4Vn; dkim-atps=neutral Received: by simark.ca (Postfix, from userid 112) id 1D9061E067; Tue, 12 May 2026 13:34:53 -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 4D3BC1E067 for ; Tue, 12 May 2026 13:34:52 -0400 (EDT) Received: from vm01.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 665DD4BA799F for ; Tue, 12 May 2026 17:34:51 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 665DD4BA799F 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=h7c/F4Vn Received: from BYAPR05CU005.outbound.protection.outlook.com (mail-westusazlp170100001.outbound.protection.outlook.com [IPv6:2a01:111:f403:c000::1]) by sourceware.org (Postfix) with ESMTPS id 5182F4BA2E31 for ; Tue, 12 May 2026 17:34:13 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 5182F4BA2E31 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 5182F4BA2E31 Authentication-Results: sourceware.org; arc=pass smtp.remote-ip=2a01:111:f403:c000::1 ARC-Seal: i=2; a=rsa-sha256; d=sourceware.org; s=key; t=1778607253; cv=pass; b=r6oM1AxjdweIyZZ9Mndde03WShSNktRo9URZY719KbPqTlQplMgPlUnsHI/LWVqpZ5ojraGEYFU9TR+Wad6aUq496lrF5srtCjnWpB4Cz9CHtS+LaApNdpDRaRlOjd6RvDQaWtM2uRLBn+HOgqUqEwHE296gRCvgL01b09nliOE= ARC-Message-Signature: i=2; a=rsa-sha256; d=sourceware.org; s=key; t=1778607253; c=relaxed/simple; bh=9Tnz9u+nH3gXk4CKVOl75554kBnhn48tcuV9KZUb4rM=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=JjJVCplqk/3MOZqaN47sy13dpef9Gx8BYxGS7mYPtn20E3/TguN1eKtB9J+a4ZJyDLP9YEa9pP9LXgygwwXJ3sRWid17OD3/YCsPyBHRqk+XwxIqb194Ax19jpULeYrHPWNebJa+kYldO/rLnLZ3wBk9mltC/vbxerRPGXVMv1c= 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=h7c/F4Vn DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 5182F4BA2E31 ARC-Seal: i=1; a=rsa-sha256; s=arcselector10001; d=microsoft.com; cv=none; b=urYOvOri59LCJ9HTDfk639diwIr/wKtAnb0vciO5evxwqgp9czAQR8JRodGiHv+k0F3u048bOGnft0D6q4A7hIuFp4QGrotu2oqQ3XiUc/vLHFYzeLTA0AdD89kgtDeW7TNInfleeoh/SW5J3O12iUn8CQ6iAlkEONMD0asuOqw7sLK48hZR4GzuGvN1YPsGmX5VG9fHEaPte2NONpizb5JyoTcW06aAIVKo/Xes8j6ymm3NeAXnYeHxCqCau7mq+kAXD8MpEesK/nz+hvmL6uwDnBEYF54sUKnOqsn9aZQHo4fMTZIJcVq6PJP5IWCBZYJLx1xLvZG9FFvcdV/oKg== 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=kbhIYxyKOpURstpzOlJf0kCT5zFfnsshummMLTq0rnc=; b=czlSDG5fXa5FuMLqEtCe7jxdq+8L8WG0DKULLFlf5ozxCmxGNHmN0TCu7UrACmDxRwlDgXe93gFOuTskRM8oi0RNbb+VWvjrPhxPQpjV3n3asT88wIAOyIblZ3jJbagWtshHmQt7LgN+GJukKIkscn285vCakKtVOrThhl0Y5UwBW7Sb21ccf6Dtmcnv68m4IgQgt2PbJWkUYi3E1+7UOmE+te0YEMRm6XdECtXWz5eLeJTmD1O4oHc/Pq8H5MmRr1VCDKmTNaD0qaFksR6yEGsr6WkhEgQqKw+mi/0jG0iDi/d1MiobDyr3M5QRvazdQCEzTvJlSdIIgVmtpf1iQQ== 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=kbhIYxyKOpURstpzOlJf0kCT5zFfnsshummMLTq0rnc=; b=h7c/F4VnpMVymlmCTbsIxqiXIeHilFQ7aZ2ScOwvs62f7Diwfpg+ssViisVDZu/HiLHa56vJIxlV9e3d1hFB+ZY6VNNNuMTE57jE0NWQlpzSS33O1RJYpTVEmVXavtydzXBEg18/IqAUgy+FpSgIxCMSKVczCvpE+bMOxc2pHf8= Received: from SA0PR13CA0028.namprd13.prod.outlook.com (2603:10b6:806:130::33) by IA0PR12MB8895.namprd12.prod.outlook.com (2603:10b6:208:491::5) 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 17:34:04 +0000 Received: from SA2PEPF000015C8.namprd03.prod.outlook.com (2603:10b6:806:130:cafe::1) by SA0PR13CA0028.outlook.office365.com (2603:10b6:806:130::33) with Microsoft SMTP Server (version=TLS1_3, cipher=TLS_AES_256_GCM_SHA384) id 15.21.25.17 via Frontend Transport; Tue, 12 May 2026 17:34:03 +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 SA2PEPF000015C8.mail.protection.outlook.com (10.167.241.198) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_256_GCM_SHA384) id 15.21.25.13 via Frontend Transport; Tue, 12 May 2026 17:34:03 +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; Tue, 12 May 2026 12:34:03 -0500 From: Tankut Baris Aktemur To: CC: , Subject: [PATCH v2] gdb/amd-dbgapi-target: implement the 'interrupt' target op Date: Tue, 12 May 2026 12:33:18 -0500 Message-ID: <20260512173318.2727821-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: SA2PEPF000015C8:EE_|IA0PR12MB8895:EE_ X-MS-Office365-Filtering-Correlation-Id: e01c1601-0d0a-4c5c-2be9-08deb04ca951 X-MS-Exchange-SenderADCheck: 1 X-MS-Exchange-AntiSpam-Relay: 0 X-Microsoft-Antispam: BCL:0; ARA:13230040|1800799024|36860700016|82310400026|376014|13003099007|56012099003|11063799003|18002099003|3023799003; X-Microsoft-Antispam-Message-Info: PBHhkMEJRjuNcIE/17NHeT7Pw2c8E1LxiALe4UO32r+QCKjJ8cKtZWtrjV7e5J5QXZghhO1027N/GIzwX/SOqbcGBarky0BB/ekwVLBFb7GUGPVHgPbobXSqcOMMvE7NzZm6h16EkoWRnL91x9oqVYI2lG+wlpQWwGms8hOkNvpZvIoxyZZHdPUn4QSYTzPVQGo6mJsjTeyi42QtGHO7hLTBSbULl4i+jbpIo2Cc5i+jp7XJLnlT8qeNMXDgsH1PMq3yrbvDQAb9fljsnVwXKbQD58fKFPltbpZ84FAwuawU0I/6ivdLgjYhMfuS48Qx01zb4fj4+wk+rF5RINTjhX4Y0eqW5UTzwd0It3TZBeYDHmrllg227HPZnnIrW2XLzOH9nthuu7fhvjTsLKW+coIGUatUxoZAUfNb1/f2E9w+eDPM9TNNpGf9zG87is9lzR4Ox14jGCmznRiQY5tWEvtACEUiJfjVKKgD1oMhjVoiwf7kC0hNOWMxXYhcFGHdl/L9UiBHkxmqGdGoNdFl0dFLRUIIYyHv7kZgsQKrimnZnRrK3GBRl0iEHty0q9eIn1talucOxHHeN4gl2tPncFjNvCAaEVeHWdw/NCsNsIMefDZyc+giTUj8Y7SjTEBxyCXaeOCmwPDoAuW0dbPy3p43AlCe3sJMYUEM2EdYD7SMMjUJF6aDXq2AcNcU3qTuoQlAq1ZDX6C0pb9oxELlhQniDFVcB8DeAyVkf6GKfvY= 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)(1800799024)(36860700016)(82310400026)(376014)(13003099007)(56012099003)(11063799003)(18002099003)(3023799003); DIR:OUT; SFP:1101; X-MS-Exchange-AntiSpam-MessageData-ChunkCount: 1 X-MS-Exchange-AntiSpam-MessageData-0: 8CRk29ejUf7E1skUe48oUxXNvUzGykrfbzWJK/VqLDCexigRPmxj6sJIkRSVjarelSX5iN5SqsUJSP0xvHeQKFPJlcneaQ8C573F/wv4guRbFxKZSLLsyNS9WSuuYFDd02qZnqNmM/xCPocVrqa04q+VgiBT3KRwYqVtZUzqfFcYB3Kh/lpTgPF+Y2U4qPHBSw+a53unaibjKD1l2SSYZi4erI6/ikgoCQu4pYj3ScGCrVw7tskQmYfcK1bOZ68apR9fFZlQHd/LLV+1FU5Wpy/fPZPYNtHUd1CqGcPAvhmGKMYW/J0WZ7GWZp+Vh0nVGkXSqJtaKvMQRJuRrkxiE1x6cy6VLrW1kxjzyKnrMzxHUqsB5/9CXsJ7B9CSTwxU78ZgMRj8/De+D805jcBTC/iHFViFKCrBr9Q93IH1hDVXEzv+jHcOUX4PO+z/BpTw X-OriginatorOrg: amd.com X-MS-Exchange-CrossTenant-OriginalArrivalTime: 12 May 2026 17:34:03.7202 (UTC) X-MS-Exchange-CrossTenant-Network-Message-Id: e01c1601-0d0a-4c5c-2be9-08deb04ca951 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: SA2PEPF000015C8.namprd03.prod.outlook.com X-MS-Exchange-CrossTenant-AuthAs: Anonymous X-MS-Exchange-CrossTenant-FromEntityHeader: HybridOnPrem X-MS-Exchange-Transport-CrossTenantHeadersStamped: IA0PR12MB8895 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. Approved-by: Lancelot Six (amdgpu) --- gdb/amd-dbgapi-target.c | 7 +++ gdb/doc/gdb.texinfo | 4 +- gdb/testsuite/gdb.rocm/interrupt-single.cpp | 43 ++++++++++++++++++ gdb/testsuite/gdb.rocm/interrupt-single.exp | 48 +++++++++++++++++++++ 4 files changed, 101 insertions(+), 1 deletion(-) 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/doc/gdb.texinfo b/gdb/doc/gdb.texinfo index f6b3b114278..d4d8a20b5bb 100644 --- a/gdb/doc/gdb.texinfo +++ b/gdb/doc/gdb.texinfo @@ -27872,7 +27872,9 @@ If no CPU thread is running, then @samp{Ctrl-C} is not able to stop @acronym{AMD GPU} threads. This can happen for example if you enable @code{scheduler-locking} after the whole program stopped, and then resume an @acronym{AMD GPU} thread. The only way to unblock the situation is to kill the -@value{GDBN} process. +@value{GDBN} process. Alternatively you can resume the @acronym{AMD GPU} +thread in the background using @code{continue&} (@pxref{Background Execution}) +and then use @code{interrupt} to stop the thread. @anchor{AMD GPU Attaching Restrictions} @item diff --git a/gdb/testsuite/gdb.rocm/interrupt-single.cpp b/gdb/testsuite/gdb.rocm/interrupt-single.cpp new file mode 100644 index 00000000000..fc8d2cca697 --- /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, 1>>> (); + 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..23ae02a0c74 --- /dev/null +++ b/gdb/testsuite/gdb.rocm/interrupt-single.exp @@ -0,0 +1,48 @@ +# 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 "loop" {allow-pending}]} { + return + } + + # 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