From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from simark.ca by simark.ca with LMTP id QPdBBFv0QmhxRAMAWB0awg (envelope-from ) for ; Fri, 06 Jun 2025 09:59:55 -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=BDy/57vZ; dkim-atps=neutral Received: by simark.ca (Postfix, from userid 112) id 0C06D1E11C; Fri, 6 Jun 2025 09:59:55 -0400 (EDT) X-Spam-Checker-Version: SpamAssassin 4.0.1 (2024-03-25) on simark.ca X-Spam-Level: X-Spam-Status: No, score=-10.1 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, RCVD_IN_VALIDITY_RPBL,RCVD_IN_VALIDITY_SAFE autolearn=ham autolearn_force=no version=4.0.1 Received: from server2.sourceware.org (server2.sourceware.org [8.43.85.97]) (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 E19571E089 for ; Fri, 6 Jun 2025 09:59:53 -0400 (EDT) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 62B0A385AC33 for ; Fri, 6 Jun 2025 13:59:53 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 62B0A385AC33 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=BDy/57vZ Received: from NAM12-DM6-obe.outbound.protection.outlook.com (mail-dm6nam12on2061e.outbound.protection.outlook.com [IPv6:2a01:111:f403:2417::61e]) by sourceware.org (Postfix) with ESMTPS id 9BE7D3858C50 for ; Fri, 6 Jun 2025 13:59:17 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 9BE7D3858C50 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 9BE7D3858C50 Authentication-Results: server2.sourceware.org; arc=pass smtp.remote-ip=2a01:111:f403:2417::61e ARC-Seal: i=2; a=rsa-sha256; d=sourceware.org; s=key; t=1749218357; cv=pass; b=w1/Vt9dfzziU9eiGb0Z7egzbBJkzD5ywX4JKTu04y5KuCjnXhK1CDEXDacVYOiUB840mhTF1oaQcxJibyFyy/kF6TmEvKrv/+6TKpumd8B7mHZBL5gMGMsj3PyP1j+Vy6QVQUQAMKN8mg5Y3VtIcIsWKx9PnHhT1SCiwOn5nA3U= ARC-Message-Signature: i=2; a=rsa-sha256; d=sourceware.org; s=key; t=1749218357; c=relaxed/simple; bh=y7KCWRV25AJO7JyfFg619omfG+IzDLPs4N3gG+u7amo=; h=DKIM-Signature:Date:From:To:Subject:Message-ID:MIME-Version; b=B2TJO+4ki7MUTC1TtEXkQx6FNwZFkg/horBclzyE+kPV3hltVBod7d13PZbKGSB763GHR5izQfL44aknmsxhjbvAUz9i9EQF30RzszteCYfCyS9HPiiNMyBAEjrj9mPyKgdHG9TAH4lop973qJzKqXvdIclkjVAwqQjZc7jrTrs= ARC-Authentication-Results: i=2; server2.sourceware.org DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 9BE7D3858C50 ARC-Seal: i=1; a=rsa-sha256; s=arcselector10001; d=microsoft.com; cv=none; b=KBfl3Mz17AyKH25maqjeicAn8drcwFNQEXgxdm6oXSZzrH2HWs9S7h5DxZlmgWNTbCtULXs+LTK/vagfG2923obGyPTHkQt+jXeu+rkezwMrOWTPzzpEWLgmJf6Cj39sJRxHFSL6VrQoVGova5uW73S0iuh4DCBV7Xg35YTxGRW1XVfMcYgl1vEKlzZ2H5cxfjjpTIbcW1s3U+Ao+j/RxfjJMvWKCQwJVcShoS2R+xjRvZQrECaaZqxsOApysEFn73Utn4hv+Mtz8f4rs9f38Mg8xJbwt4tI74ZFq1HS489vYxmqY+e8ZQbqiAdW1VbdYbUxvFQ1RnmYP/HeEVowPw== 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=RPSHzOMV9c+PXpoRbLOsMLj9YdXNI4mnPWbkNLUOcZw=; b=D3wAqpUxSt8K9TbohO/8dHQ04R7el17GgH+EyfesoyaMbafyeEFCMNEtZMyUDZqQt1UQsHfo2A5gaA0kJ4c4m5Bbpc9Nw+9BZuQ0k3seuv72hf9Fb2ESEYIEOpKhItiRVsEcr1kgsVZqxVNv6kUzZaqfIn2ddQkfInVniMoU5pmnjs4612iikgPRCuTPWOLMVQvUYExMVue5iEkLXdGyrDjry1fWkf/H6jUJktJNh+aLxOpFW2vv8Pog+wNx5F1QBAxI/zqVAsXhr1gARbxWUg2a4s/4W+ff7MZpO5dhH7qvG4S1lsfmV1VvvM+JZ7mG0XaHGKY/ITq6mPmU4pCP7w== ARC-Authentication-Results: i=1; mx.microsoft.com 1; spf=pass (sender ip is 165.204.84.17) smtp.rcpttodomain=efficios.com 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=RPSHzOMV9c+PXpoRbLOsMLj9YdXNI4mnPWbkNLUOcZw=; b=BDy/57vZB/2u0jcN5kKqKp54nWSU7p64L28WYl1jA0zVnaN6VB1KzLjsDzLYbP3K+XWtM7YA+DZPMHVG85UKQq9YJ1wHUXqUtUDkOpbIMenaZdfzCcCq1p1Z9PHIN3E0blApXytHkiH8UYvCnPzceIbMdzyMkj3QNTLV+JI92Bg= Received: from BY5PR17CA0071.namprd17.prod.outlook.com (2603:10b6:a03:167::48) by CH2PR12MB4214.namprd12.prod.outlook.com (2603:10b6:610:aa::21) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_256_GCM_SHA384) id 15.20.8792.34; Fri, 6 Jun 2025 13:59:12 +0000 Received: from MWH0EPF000989E8.namprd02.prod.outlook.com (2603:10b6:a03:167:cafe::3b) by BY5PR17CA0071.outlook.office365.com (2603:10b6:a03:167::48) with Microsoft SMTP Server (version=TLS1_3, cipher=TLS_AES_256_GCM_SHA384) id 15.20.8746.31 via Frontend Transport; Fri, 6 Jun 2025 13:59:12 +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=SATLEXMB04.amd.com; pr=C Received: from SATLEXMB04.amd.com (165.204.84.17) by MWH0EPF000989E8.mail.protection.outlook.com (10.167.241.135) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_128_GCM_SHA256) id 15.20.8792.29 via Frontend Transport; Fri, 6 Jun 2025 13:59:11 +0000 Received: from khazad-dum (10.180.168.240) by SATLEXMB04.amd.com (10.181.40.145) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_128_GCM_SHA256) id 15.1.2507.39; Fri, 6 Jun 2025 08:59:08 -0500 Date: Fri, 6 Jun 2025 14:58:58 +0100 From: Lancelot SIX To: Simon Marchi CC: Subject: Re: [PATCH 5/5] gdb/amd-dbgapi: disable forward progress requirement in amd_dbgapi_target_breakpoint::check_status Message-ID: <666t5msvuyylv2yioapk56jqvjbsxsc7spkf473uu6pfguytfe@hvn54ws4ctzp> References: <20250605201657.418206-5-simon.marchi@efficios.com> MIME-Version: 1.0 Content-Type: text/plain; charset="us-ascii" Content-Disposition: inline In-Reply-To: <20250605201657.418206-5-simon.marchi@efficios.com> X-Originating-IP: [10.180.168.240] X-ClientProxiedBy: SATLEXMB04.amd.com (10.181.40.145) To SATLEXMB04.amd.com (10.181.40.145) X-EOPAttributedMessage: 0 X-MS-PublicTrafficType: Email X-MS-TrafficTypeDiagnostic: MWH0EPF000989E8:EE_|CH2PR12MB4214:EE_ X-MS-Office365-Filtering-Correlation-Id: db01a5e7-f294-4007-dd7b-08dda50250a8 X-MS-Exchange-SenderADCheck: 1 X-MS-Exchange-AntiSpam-Relay: 0 X-Microsoft-Antispam: BCL:0; ARA:13230040|36860700013|376014|82310400026|1800799024; X-Microsoft-Antispam-Message-Info: =?us-ascii?Q?3m7tPqPIBNNCPoTABi1OOowJnXfmZc9Nlzbk5Lg5vUdAois00hGk5qW/iST6?= =?us-ascii?Q?8JYzsHYQ6/uNpD5ABJ0ft8wQtBc1qND1I2ad/EhBGENUhFot78KZE9X/z6Ze?= =?us-ascii?Q?PFqukScDqrpt3T/7pQecHlbrsqmmSsEVJlnsd6tPJrRbPgwoi0wBW5EGaEpV?= =?us-ascii?Q?GyRfUmCjIxKSl9AoimjipDOQAR29azDeDLBupkUEwlIQzxUYW97TsbVlkLsO?= =?us-ascii?Q?SxY1/GPu8q+4oF6JSbgB9UD75/7QJrv3HxL7yyp/pZhkzNSCps2/mcUn5CT8?= =?us-ascii?Q?csxuPVSLs8azhvp03sY2CYaIKIUqjI0vdb7UmuqO+21ThuyC6e9eFTpbVWGA?= =?us-ascii?Q?9ZdowqhVrgJ8UL2RCcP8kKXig3XAgMJJQGGFi9n3Sy9RQzcCr56QCS76Kf//?= =?us-ascii?Q?wmIlr9TLX4kP2LFURmAPLvFNXqf4ZkkJaykOPRwbJ5pn/HZxyTMZR+i3SLH0?= =?us-ascii?Q?7DydOylQFoq8+b7zp5VCPj7HTF+BmBibs3/0OA+qXxYZrrZGjCmrZnu2xnpf?= =?us-ascii?Q?Wf7+61g4/ydEz5hPEAR9IzFGUfKYqeUJpcw8nAI8eZDWbKC57tTw5iZ047K3?= =?us-ascii?Q?LszY2io8ZKLI8YaMl7dBGsMqw7dBMec6fiaPEr4ZVoY++7NOdg6bZTm8D+me?= =?us-ascii?Q?m3mq67WOia+R7crLlt/ZeUENdeF9N/bWeh3H6afsuqJG1B0DJ2nLv4qcu8Uc?= =?us-ascii?Q?z2p5B+6Bkb+auXeoIs63d7ay1U+O6Yz6bG6z1ho0cYXY318pSogrxeC2dFLR?= =?us-ascii?Q?siQaRreV8gWMXh2ADXJpOZa8rIzocvVdQHYfyhSybGER7Z9x0yofi+tiatAO?= =?us-ascii?Q?3DyqiOwm1PDnaDBwAclZsrnYWb5emn3pfH2lHzMe5Dx53nDEMsNZJR+bSAQV?= =?us-ascii?Q?zC9bN+ywN51aZ7c61yDjcM+ASd3lLTHnKMap+eeV0NBtmdA4dljhGCkg/zaI?= =?us-ascii?Q?a6atFAOA9JJ/JKVXJC7+6WbQN5fbjF3XnsYQ+buUpshrRVgO7khLvlZINUJB?= =?us-ascii?Q?sWfZsx/N2o6Wuy7cjMFMxS+3KKXLmwBmUaUrAM5hsB2jMU+JYBdhWRmUSWyG?= =?us-ascii?Q?rlRaBqkKlKHCLZrneHogWFul9saMAA9CwyisAtx2DV0gj6to9nOAmyy1oMc1?= =?us-ascii?Q?nNQXp4GfY3DtDIO67WmcPTtS+Y0+PO7qr0r87yRjSbEfewKSSAg5tKJfjkOQ?= =?us-ascii?Q?Vl/B6V55hQP1YlisfMIRPdifXkjzURYRqtJ3fJ3Y7oIpmm3+YlpXfgJXsxHN?= =?us-ascii?Q?P066egYYkvtUGWU1d7V3Etn1tDq67AQr965n6+XvP/v/bgTF4pKNppt3eCvs?= =?us-ascii?Q?5GgCz/aUQxACuujRBwUAmm4Lh4RxjVjmog7yKM7ls1aH/COUqyDIAk/+xZMz?= =?us-ascii?Q?+yr3HTS59bJF6iwoEvh8DMMtx1MpNfsaa3OqqjAjzm3NQP/3cHvut22D258s?= =?us-ascii?Q?+Z7WuWQe53ZXO/dj82oy48sfHCiCTLlGU5x+yZB/++z9TJvDK2E0VppbY8Ez?= =?us-ascii?Q?6rJ/VLFZZJg0TdBM2lUTNbkTKbK/wA+K2DzrKDmHLfHt4ftwf3l9VEA0YQ?= =?us-ascii?Q?=3D=3D?= X-Forefront-Antispam-Report: CIP:165.204.84.17; CTRY:US; LANG:en; SCL:1; SRV:; IPV:CAL; SFV:NSPM; H:SATLEXMB04.amd.com; PTR:InfoDomainNonexistent; CAT:NONE; SFS:(13230040)(36860700013)(376014)(82310400026)(1800799024); DIR:OUT; SFP:1101; X-OriginatorOrg: amd.com X-MS-Exchange-CrossTenant-OriginalArrivalTime: 06 Jun 2025 13:59:11.6769 (UTC) X-MS-Exchange-CrossTenant-Network-Message-Id: db01a5e7-f294-4007-dd7b-08dda50250a8 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=[SATLEXMB04.amd.com] X-MS-Exchange-CrossTenant-AuthSource: MWH0EPF000989E8.namprd02.prod.outlook.com X-MS-Exchange-CrossTenant-AuthAs: Anonymous X-MS-Exchange-CrossTenant-FromEntityHeader: HybridOnPrem X-MS-Exchange-Transport-CrossTenantHeadersStamped: CH2PR12MB4214 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 Thu, Jun 05, 2025 at 04:16:28PM -0400, Simon Marchi wrote: > ROCgdb handles target events very slowly when running a test case like > this, where a breakpoint is preset on HipTest::vectorADD: > > for (int i=0; i < numDevices; ++i) { > HIPCHECK(hipSetDevice(i)); > hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, stream[i], > static_cast(A_d[i]), static_cast(B_d[i]), C_d[i], N); > } > > What happens is: > > - A kernel is launched > - The internal runtime breakpoint is hit during the second > hipLaunchKernelGGL call, which causes > amd_dbgapi_target_breakpoint::check_status to be called > - Meanwhile, all waves of the kernel hit the breakpoint on vectorADD > - amd_dbgapi_target_breakpoint::check_status calls process_event_queue, > which pulls the thousand of breakpoint hit events from the kernel > - As part of handling the breakpoint hit events, we write the PC of the > waves that stopped to decrement it. Because the forward progress > requirement is not disabled, this causes a suspend/resume of the > queue each time, which is time-consuming. > > The stack trace where this all happens is: > > #32 0x00007ffff6b9abda in amd_dbgapi_write_register (wave_id=..., register_id=..., offset=0, value_size=8, value=0x7fffea9fdcc0) at /home/smarchi/src/amd-dbgapi/src/register.cpp:587 > #33 0x00005555588c0bed in amd_dbgapi_target::store_registers (this=0x55555c7b1d20 , regcache=0x507000002240, regno=470) at /home/smarchi/src/wt/amd/gdb/amd-dbgapi-target.c:2504 > #34 0x000055555a5186a1 in target_store_registers (regcache=0x507000002240, regno=470) at /home/smarchi/src/wt/amd/gdb/target.c:3973 > #35 0x0000555559fab831 in regcache::raw_write (this=0x507000002240, regnum=470, src=...) at /home/smarchi/src/wt/amd/gdb/regcache.c:890 > #36 0x0000555559fabd2b in regcache::cooked_write (this=0x507000002240, regnum=470, src=...) at /home/smarchi/src/wt/amd/gdb/regcache.c:915 > #37 0x0000555559fc3ca5 in regcache::cooked_write (this=0x507000002240, regnum=470, val=140737323456768) at /home/smarchi/src/wt/amd/gdb/regcache.c:850 > #38 0x0000555559fab09a in regcache_cooked_write_unsigned (regcache=0x507000002240, regnum=470, val=140737323456768) at /home/smarchi/src/wt/amd/gdb/regcache.c:858 > #39 0x0000555559fb0678 in regcache_write_pc (regcache=0x507000002240, pc=0x7ffff62bd900) at /home/smarchi/src/wt/amd/gdb/regcache.c:1460 > #40 0x00005555588bb37d in process_one_event (event_id=..., event_kind=AMD_DBGAPI_EVENT_KIND_WAVE_STOP) at /home/smarchi/src/wt/amd/gdb/amd-dbgapi-target.c:1873 > #41 0x00005555588bbf7b in process_event_queue (process_id=..., until_event_kind=AMD_DBGAPI_EVENT_KIND_BREAKPOINT_RESUME) at /home/smarchi/src/wt/amd/gdb/amd-dbgapi-target.c:2006 > #42 0x00005555588b1aca in amd_dbgapi_target_breakpoint::check_status (this=0x511000140900, bs=0x50600014ed00) at /home/smarchi/src/wt/amd/gdb/amd-dbgapi-target.c:890 > #43 0x0000555558c50080 in bpstat_stop_status (aspace=0x5070000061b0, bp_addr=0x7fffed0b9ab0, thread=0x518000026c80, ws=..., stop_chain=0x50600014ed00) at /home/smarchi/src/wt/amd/gdb/breakpoint.c:6126 > #44 0x000055555984f4ff in handle_signal_stop (ecs=0x7fffeaa40ef0) at /home/smarchi/src/wt/amd/gdb/infrun.c:7169 > #45 0x000055555984b889 in handle_inferior_event (ecs=0x7fffeaa40ef0) at /home/smarchi/src/wt/amd/gdb/infrun.c:6621 > #46 0x000055555983eab6 in fetch_inferior_event () at /home/smarchi/src/wt/amd/gdb/infrun.c:4750 > #47 0x00005555597caa5f in inferior_event_handler (event_type=INF_REG_EVENT) at /home/smarchi/src/wt/amd/gdb/inf-loop.c:42 > #48 0x00005555588b838e in handle_target_event (client_data=0x0) at /home/smarchi/src/wt/amd/gdb/amd-dbgapi-target.c:1513 > > Fix that performance problem by disabling the forward progress > requirement in amd_dbgapi_target_breakpoint::check_status, before > calling process_event_queue, so that we can process all events > efficiently. > > Since the same performance problem could theoritically happen any time > process_event_queue is called with forward progress requirement enabled, > add an assert to ensure that forward progress requirement is disabled > when process_event_queue is invoked. This makes it necessary to add a > require_forward_progress call to amd_dbgapi_finalize_core_attach. It > looks a bit strange, since core files don't have execution, but it > doesn't hurt. > > Add a test that replicates this scenario. The test launches a kernel > that hits a breakpoint (with an always false condition) repeatedly. > Meanwhile, the host process loads an unloads a code object, causing > check_status to be called. > > Bug: SWDEV-482511 > Change-Id: Ida86340d679e6bd8462712953458c07ba3fd49ec Hi Simon, This looks overall good to me. > --- > gdb/amd-dbgapi-target.c | 6 ++ > .../code-object-load-while-breakpoint-hit.cpp | 70 +++++++++++++++++++ > .../code-object-load-while-breakpoint-hit.exp | 68 ++++++++++++++++++ > 3 files changed, 144 insertions(+) > create mode 100644 gdb/testsuite/gdb.rocm/code-object-load-while-breakpoint-hit.cpp > create mode 100644 gdb/testsuite/gdb.rocm/code-object-load-while-breakpoint-hit.exp > > diff --git a/gdb/amd-dbgapi-target.c b/gdb/amd-dbgapi-target.c > index e86d7f34d0f6..44913bbd7532 100644 > --- a/gdb/amd-dbgapi-target.c > +++ b/gdb/amd-dbgapi-target.c > @@ -568,6 +568,8 @@ amd_dbgapi_target_breakpoint::check_status (struct bpstat *bs) > if (action == AMD_DBGAPI_BREAKPOINT_ACTION_RESUME) > return; > > + require_forward_progress (*info, false); > + > /* If the action is AMD_DBGAPI_BREAKPOINT_ACTION_HALT, we need to wait until > a breakpoint resume event for this breakpoint_id is seen. */ > amd_dbgapi_event_id_t resume_event_id > @@ -1352,6 +1354,10 @@ static amd_dbgapi_event_id_t > process_event_queue (amd_dbgapi_inferior_info &info, > amd_dbgapi_event_kind_t until_event_kind) > { > + /* Pulling events with forward progress required may result in bad > + performance, make sure it is not required. */ > + gdb_assert (!info.forward_progress_required); > + > while (true) > { > amd_dbgapi_event_id_t event_id; > diff --git a/gdb/testsuite/gdb.rocm/code-object-load-while-breakpoint-hit.cpp b/gdb/testsuite/gdb.rocm/code-object-load-while-breakpoint-hit.cpp > new file mode 100644 > index 000000000000..e97809335569 > --- /dev/null > +++ b/gdb/testsuite/gdb.rocm/code-object-load-while-breakpoint-hit.cpp > @@ -0,0 +1,70 @@ I think you are missing the copyright header for this file. > +#ifdef DEVICE > + > +#include > + > +constexpr unsigned int NUM_BREAKPOINT_HITS = 5; > + > +static __device__ void > +break_here () > +{ > + asm ("s_nop 0"); I don't expect this inline asssembly is necessary. > +} > + > +extern "C" __global__ void > +kernel () > +{ > + for (int n = 0; n < NUM_BREAKPOINT_HITS; ++n) > + break_here (); > +} > + > +#else > + > +#include > +#include > + > +constexpr unsigned int NUM_ITEMS_PER_BLOCK = 256; > +constexpr unsigned int NUM_BLOCKS = 128; > +constexpr unsigned int NUM_ITEMS = NUM_ITEMS_PER_BLOCK * NUM_BLOCKS; > +constexpr unsigned int NUM_LOAD_UNLOADS = 5; > + > +#define CHECK(cmd) \ > + { \ > + hipError_t error = cmd; \ > + if (error != hipSuccess) \ > + { \ > + fprintf (stderr, "error: '%s'(%d) at %s:%d\n", \ > + hipGetErrorString (error), error, __FILE__, __LINE__); \ > + exit (EXIT_FAILURE); \ > + } \ > + } > + > +int > +main (int argc, const char **argv) > +{ > + if (argc != 2) > + { > + fprintf (stderr, "Usage: %s \n", argv[0]); > + return 1; > + } > + > + const auto module_path = argv[1]; > + hipModule_t module; > + CHECK (hipModuleLoad (&module, module_path)); > + > + /* Launch the kernel. */ > + hipFunction_t function; > + CHECK (hipModuleGetFunction (&function, module, "kernel")); > + CHECK (hipModuleLaunchKernel (function, NUM_BLOCKS, 1, 1, > + NUM_ITEMS_PER_BLOCK, 1, 1, 0, nullptr, nullptr, > + nullptr)); > + > + /* Load and unload the module many times. */ > + for (int i = 0; i < NUM_LOAD_UNLOADS; ++i) > + { > + hipModule_t dummy_module; > + CHECK (hipModuleLoad (&dummy_module, module_path)); > + CHECK (hipModuleUnload (dummy_module)); > + } > +} > + > +#endif > diff --git a/gdb/testsuite/gdb.rocm/code-object-load-while-breakpoint-hit.exp b/gdb/testsuite/gdb.rocm/code-object-load-while-breakpoint-hit.exp > new file mode 100644 > index 000000000000..b8d003f37032 > --- /dev/null > +++ b/gdb/testsuite/gdb.rocm/code-object-load-while-breakpoint-hit.exp > @@ -0,0 +1,68 @@ > +# Copyright (C) 2025 Advanced Micro Devices, Inc. All rights reserved. The copyright should be to the FSF, not AMD. Given the elements above are fixed, Approved-by: Lancelot Six (amdgpu) Best, Lanecelot. > + > +# 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 . > + > +# This test verifies what happens when a code object list update happens at the > +# same time as some wave stop events are reported. It was added following a > +# performance bug fix, where forward progress requirement disabled when > +# pulling events from amd-dbgapi in amd_dbgapi_target_breakpoint::check_status. > +# > +# The test launches a kernel that hits a breakpoint with an always false > +# condition a certain number of times. Meanwhile, the host loads and unloads > +# a code object in a loop, causing check_status to be called. The hope is that > +# check_status, when calling process_event_queue, will pull many WAVE_STOP > +# events from the kernel hitting the breakpoint. > +# > +# Without the appropriate fix (of disabling forward progress requirement in > +# check_status), GDB would hit the newly-added assert in process_event_queue, > +# which verifies that forward progress requirement is disabled. Even without > +# this assert, the test would likely time out (depending on the actual timeout > +# value). > + > +load_lib rocm.exp > +standard_testfile .cpp > +require allow_hipcc_tests > + > +# Build the host executable. > +if { [build_executable "failed to prepare" \ > + $testfile $srcfile {debug hip}] == -1 } { > + return -1 > +} > + > +set hipmodule_path [standard_output_file ${testfile}.co] > + > +# Build the kernel object file. > +if { [gdb_compile $srcdir/$subdir/$srcfile \ > + $hipmodule_path object \ > + { debug hip additional_flags=--genco additional_flags=-DDEVICE } ] != "" } { > + return -1 > +} > + > +proc do_test { } { > + with_rocm_gpu_lock { > + clean_restart $::binfile > + gdb_test_no_output "set args $::hipmodule_path" "set args" > + > + if { ![runto_main] } { > + return > + } > + > + gdb_test "with breakpoint pending on -- break break_here if 0" > + gdb_continue_to_end "continue to end" "continue" 1 > + } > +} > + > +do_test > -- > 2.49.0