From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from simark.ca by simark.ca with LMTP id NkswM1atg2n5rygAWB0awg (envelope-from ) for ; Wed, 04 Feb 2026 15:34:30 -0500 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=jdew1bUU; dkim-atps=neutral Received: by simark.ca (Postfix, from userid 112) id BD69D378002; Wed, 04 Feb 2026 15:34:30 -0500 (EST) X-Spam-Checker-Version: SpamAssassin 4.0.1 (2024-03-25) on simark.ca X-Spam-Level: X-Spam-Status: No, score=-2.6 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,RDNS_NONE autolearn=ham autolearn_force=no version=4.0.1 Received: from vm01.sourceware.org (unknown [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 E0B5B1E08D for ; Wed, 04 Feb 2026 15:34:26 -0500 (EST) Received: from vm01.sourceware.org (localhost [127.0.0.1]) by sourceware.org (Postfix) with ESMTP id 5D7784BA2E10 for ; Wed, 4 Feb 2026 20:34:26 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 5D7784BA2E10 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=jdew1bUU Received: from DM5PR21CU001.outbound.protection.outlook.com (mail-centralusazon11011002.outbound.protection.outlook.com [52.101.62.2]) by sourceware.org (Postfix) with ESMTPS id A67C24BA2E09 for ; Wed, 4 Feb 2026 20:33:48 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org A67C24BA2E09 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 A67C24BA2E09 Authentication-Results: server2.sourceware.org; arc=pass smtp.remote-ip=52.101.62.2 ARC-Seal: i=2; a=rsa-sha256; d=sourceware.org; s=key; t=1770237228; cv=pass; b=dRjV6jIaIzYYIoEy9blQAbI1RfmYQpCxrYR/Ru4H4dC0kbYiN1c44kkskFTqyJZYpPVXFEZQKF9lwq+lX5wNVn9E/lrqUu1tKPhzMtv6hB0WzPUCtt81Hvp1pw8DpNM/Z+Z+xZVCl581gyRP9ciPzkqrZyv/daogG+SpkQoeWbY= ARC-Message-Signature: i=2; a=rsa-sha256; d=sourceware.org; s=key; t=1770237228; c=relaxed/simple; bh=89fHmfb0n2U3jMblTiye86wAtCSHPVSS+7NaGtIs6aQ=; h=DKIM-Signature:Date:From:To:Subject:Message-ID:MIME-Version; b=SZhWt6NdbbOz1mCvkf2ltGjRAwyg9Gq8hA70009Nqcm46ZHwKO/8Dn0qzSQGzg3XrIniwckIIOFB90013tfV6EscLcKgrji+UGg3fwOW5CiCBMiKR6ZZAL0QWkGYfwqKLJhHca1TuVJU6AQKyDfCAylVbKTXBPxN7PldhYMFwCg= ARC-Authentication-Results: i=2; server2.sourceware.org DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org A67C24BA2E09 ARC-Seal: i=1; a=rsa-sha256; s=arcselector10001; d=microsoft.com; cv=none; b=Hr2bTb73WoKLS/9eyI733urujdYGdsJ/OteXZy0bSgTI924TWXACePsLuxibb6hoKxkJAyL7mBx8z+3Ic4b40QlIL9MDGPFTSzO3P1jr1Kp/OJyPoBYy6vv1GkmstS6bjTz2e0+sU0nuT6PzzMbwiu7DobPgZeG9vSVl0ohlKcZG8GYyuZl38VmuL2tvhz+4FMgdRqIVdPcCzTrkr7ZTw3NptnvNz1Kqt6HHbigxY83SC/C1qMoi2FsIq4LnzXMg2DZV2RM7nx6JpTD/Zo7T22RXKXaQaSMSi8wTfSK6W0f9jQJ0n0jPFZGlneYDIkpM65SpbW1AcUVXnm73i5a0dg== 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=XReXJDMj8HQusQjUY3hBIf90GbTwg3f57FC/Q71gf4g=; b=exmU1SY7YKow/j0A+KM612Tt8T9Nyn2z0a60TTCDJNIgeNtmrRsNkFnU4ack5zIh0+Gl1vxp3IZOd2HaLBZnBZ2alO+NnU0Fp18ZZgqhEvfInnKEfnZDpdIv9YzdA7FGiWj+MicNUvU8n7QmqoVgh9XKsxiYbr+nUW8+TzioqtDCKJnZIkaFRHbgwJ7BGziNh3QQbBThAiHYv7mPJEs+zOfaEs0xJNP1Oegej2gqdRywj+7hjA0fSgZPHaTv2S0QtQm//jDTJcYPqpQg3yA2UOc+RlzuuVJyjL8jv2Xpt8N7o89pFlk1A+oKxlcOFYUh07+VC8FALJ6eMMZs+3Gepw== ARC-Authentication-Results: i=1; mx.microsoft.com 1; spf=pass (sender ip is 165.204.84.17) smtp.rcpttodomain=polymtl.ca 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=XReXJDMj8HQusQjUY3hBIf90GbTwg3f57FC/Q71gf4g=; b=jdew1bUUBiU9Ge78dxa8r5Nx18kvpr+zJq5cLUt/EUsunWYF4Qhcntlc0qTt9/q6mkGsmlZ3f8sS3pOhtGfqSGHEV7SRyYZi0uNHteh6rLZ0fdzO3xB634L8V3+CZ2eHjsBS5qg4I3ehws44k9v4Fu5tVyoB50n/dTIvL6hODV4= Received: from PH7P220CA0145.NAMP220.PROD.OUTLOOK.COM (2603:10b6:510:327::9) by DS0PR12MB7993.namprd12.prod.outlook.com (2603:10b6:8:14b::14) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_256_GCM_SHA384) id 15.20.9587.12; Wed, 4 Feb 2026 20:33:42 +0000 Received: from SN1PEPF000252A1.namprd05.prod.outlook.com (2603:10b6:510:327:cafe::be) by PH7P220CA0145.outlook.office365.com (2603:10b6:510:327::9) with Microsoft SMTP Server (version=TLS1_3, cipher=TLS_AES_256_GCM_SHA384) id 15.20.9564.16 via Frontend Transport; Wed, 4 Feb 2026 20:33:38 +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 SN1PEPF000252A1.mail.protection.outlook.com (10.167.242.8) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_256_GCM_SHA384) id 15.20.9587.10 via Frontend Transport; Wed, 4 Feb 2026 20:33:41 +0000 Received: from khazad-dum (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.17; Wed, 4 Feb 2026 14:33:40 -0600 Date: Wed, 4 Feb 2026 20:33:32 +0000 From: Lancelot SIX To: Simon Marchi CC: Subject: Re: [PATCH 2/2] gdb/amd-dbgapi: add basic watchpoint support Message-ID: References: <20260124051512.731-2-simon.marchi@polymtl.ca> MIME-Version: 1.0 Content-Type: text/plain; charset="utf-8" Content-Disposition: inline Content-Transfer-Encoding: 8bit In-Reply-To: <20260124051512.731-2-simon.marchi@polymtl.ca> 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: SN1PEPF000252A1:EE_|DS0PR12MB7993:EE_ X-MS-Office365-Filtering-Correlation-Id: c55576de-8c8f-45dc-1ada-08de642caf91 X-MS-Exchange-SenderADCheck: 1 X-MS-Exchange-AntiSpam-Relay: 0 X-Microsoft-Antispam: BCL:0; ARA:13230040|36860700013|82310400026|1800799024|376014|13003099007; X-Microsoft-Antispam-Message-Info: =?utf-8?B?K0hNMENObWlkTHVrRHNhbG1RZ2tOQWcyVVlTS20rVkZ2ZkpYeUs3VVZTbGdK?= =?utf-8?B?UlJ3NUdFd0plMnlqZ2Q5UDBMWS9yaE1LZW5IY0R4T2ExVnFnM0J0dEs0Z2hP?= =?utf-8?B?K2FjTnpBZ21vNlpVNldLbkEwc1owL095TjVkMks0RU82ei90NmVYRXQ0ZTMy?= =?utf-8?B?QmNSMFprcktldWN2eGdTMFltdGJDd2tPQlRjcys0UmdpOGtVeXpmM1RZR2pj?= =?utf-8?B?emdVR0tpQVl1TXJvN1lXYjFwemJzcHdDaTc1LytVdEFIemJKYitZRlFteW1v?= =?utf-8?B?LzgvK044bk5HTFQzeS9TOFU3aWxUUCtmV2g3MDUydlk5OEQ3eXpuQ1EwcGtx?= =?utf-8?B?WlNJNjlrSjBKaytOanlyWmRhR3JSMkk3Y0dMN1BCazVoNXVQQ2VTUFNSWG5B?= =?utf-8?B?YW5RMW5IeGluSGRUdlVOTDVnVkVIVWIvRTlXTEFJMjR3U042WWlONUZBbFVo?= =?utf-8?B?VW5wZW5xR1hZN1c4bXhHWDdEc3RYbDloWE02dVRxdDVFNWZyaFgvVTlQaDVH?= =?utf-8?B?YVZtUWpIbXc3Q01DbXlrWHp5dHIvYytjMTJGS2s0aXBZNkttMEhJUklGZmJJ?= =?utf-8?B?YzRHT05PMGtCbmxhZ25HU2JWa2xEWGJ6UHhkVEs1VG1oK3FTRFV6S3RtMnAv?= =?utf-8?B?RDVqbjdPZng1N1RGMDFLRUxwdmtDYVBNZ1B0OEZ3UTZhaTZ5Yk9zSVZlS2hR?= =?utf-8?B?dFBYSEFOdHI4VWxzSkZ5Nk9LR01lWDUrTllQSnI3eC9rTlErejgvV0h0S05P?= =?utf-8?B?enlYcXRpd1FLYzc4YlBUN0R4Wks0V25HWmNCVThyZkpMT2RWa20rZm54NW1G?= =?utf-8?B?Tjk3SzZ6Rm9QZ1ZXNFRrSGdJWjJaVjZPZnR5c2JIbEZiNitTaTZaMitER0xY?= =?utf-8?B?UEhyTG5UMUdqNXQ3NEkwUFVyTlp4ZDJlTXBLTDZWaUNUSFBzM3RJa0QwMUFu?= =?utf-8?B?K1NwUy9FYjNhUkc2WDkrWHlmODFKYkt2VnJZZ04wVEpqdlBvRUt5eFBpYXMv?= =?utf-8?B?TVBvNmpCdkJGbkF1N1AwRktXYy9NbEE5VEhRQUFXWWlMMVp6V0h2SFJNVFNR?= =?utf-8?B?ekh2UFJRUzZtR3JjOTJUTnVsY1ltWFpzVkl1cFhRTElnTXIrdlR0a2tDTHFG?= =?utf-8?B?cDY5UXQ5OEEzbHJmcHVObVBQOEJuZDRmMEU3M2RrcVU3UXE2QU5xaXFKeHVk?= =?utf-8?B?OWxUZEpQQmJPLzNFbjRIdHlXMHluMWVzNjR4ekdhZ0N0d3hXbURNQWRnNVZv?= =?utf-8?B?MlVOaWh4RlRSR0ZNd3VlcVBDemUzUTZWMnRuRDM1WlpKc0dqYVFzN08rbllT?= =?utf-8?B?dENnUkRHczlGNVVvcWFaR25VKzE4WXZKWFJjcGhGamQ4Sm84TVB1TGxyVXlr?= =?utf-8?B?NXZjSWxxNUVhZUN6RkFCT0E0U04xR1hacFFEOWlSaCtZMDY5VGMzaDRWOFFD?= =?utf-8?B?YURnVVRhMDMxa1FrSGZtRGJhaG1PZHVsMjVmbVVRSWt4dHcyTjBwajIvK1pt?= =?utf-8?B?WEg2ZzFXQTNBUjdWbkVuN1AzbExwaERoVUtFbEx0QWZua210Yk84eHhTVlFJ?= =?utf-8?B?bVkvWUx4N1p1Wk52aGF6YThoRnVQRWJTbFRiMDM4aklXdmwyaFp3TVZZRncr?= =?utf-8?B?ZUhmSnFhZ25lUUZ3bFl0Z2pKemQ3NmQvWERvL0l1em1sb3MrU2RJT2NRZVVT?= =?utf-8?B?eHc0elR2b0o1eG05K3c1SXIxZURzaks0SkVGbmEzTC91TTZJZm9oQ1pjUG42?= =?utf-8?B?dTFZKzlLcG4zanVMZ3o0UVhpUnFIaE5pbGtQblM0YUhkaUo3Q0FIRndXMDBC?= =?utf-8?B?WE5VNEgxK1JOZzRnK1EwbkkvQ0RVVXdsc2VIdkdSZGFhYTlKRkFRR2RSR1Q0?= =?utf-8?B?OWVudkRKQk9QN1ArMHBTc2I3Q3hUMGRjemsxV08wb3hPSmdBSTBuTU5MdFRX?= =?utf-8?B?NG5rNCtDRDRTbTlTU2IwT2lML3laekEyVy9Edk1qSGdETm5ZWGxQejZuOVI0?= =?utf-8?B?a2Z6MXhYVkIxcVNMMGRCd2MvN1YwVHFLMmV6Rlh0dlpEQlRJQTd5bXNPVGZi?= =?utf-8?B?VnlGZC9XbldoSENFaDllTWR4bDRWb0xqRmlERTQzdThBRGdPMEJmdU9mV25O?= =?utf-8?B?Zk1qVTR4NGhEWUFOY1ZZRHhBQjlhY0IxeU9DYlJlL0VxZGZ1YmM0WUw0dEpu?= =?utf-8?B?VDc4dWQvd2dyREVWemNtaUsxSlpCWkNHcFZKNVdkS1FZMWMxUU9ZYk9qSjNX?= =?utf-8?B?QmJ6eENoZ2hYcUNSVFhEN1c4cTFBPT0=?= 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)(36860700013)(82310400026)(1800799024)(376014)(13003099007); DIR:OUT; SFP:1101; X-MS-Exchange-AntiSpam-MessageData-ChunkCount: 1 X-MS-Exchange-AntiSpam-MessageData-0: h6HB/PryA5TPjbAsoTfiVtgzHGqmWKuC3a8DJASy6dyKkgqvvYxzWhvRySIxT6KfKaCTTlnkiEgZawu5k5+Tr3oMAoMf08+c06u6XBB+KK78pbGwI6/q0ektO7YzohVb5yRoMyIOfMZL2Am2iAdp9ym2vs+F7AC++Pd+n2RDdNTjTSWeCI0kXH0E/48FoPHWZnbiqthg75DqpN3ImanPCdluArfJleL2mvKeO1qhVgkNJvgu3C+bqprqYaWp2YsQ8336QSZukyYURlbBchuqj6Ws4Pi3IntDVeZexZt4cOLUq4CQ+Rg3fuVhJ6uGf+kVZKGq36lQ76FubRzpF20D4GB09YbwZCWxpsqTTYILAe9bHtw7AaTIGN18AZO2NroFYMLvF/mZYu0bPHtPOJNBPm0w50hQDabrpusEiQ33jG1kikVMxt0VARtTunh7RUkI X-OriginatorOrg: amd.com X-MS-Exchange-CrossTenant-OriginalArrivalTime: 04 Feb 2026 20:33:41.9633 (UTC) X-MS-Exchange-CrossTenant-Network-Message-Id: c55576de-8c8f-45dc-1ada-08de642caf91 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: SN1PEPF000252A1.namprd05.prod.outlook.com X-MS-Exchange-CrossTenant-AuthAs: Anonymous X-MS-Exchange-CrossTenant-FromEntityHeader: HybridOnPrem X-MS-Exchange-Transport-CrossTenantHeadersStamped: DS0PR12MB7993 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 Simon, On Sat, Jan 24, 2026 at 12:15:00AM -0500, Simon Marchi wrote: > Add basic watchpoint support for the amd-dbgapi target. This means > placing write watchpoints on globally addressable memory. More > complexity will come eventually to allow placing watchpoints on the > various other address spaces, but that will require adding proper > support for non-default address spaces first. > > Implementation > -------------- > > I think the implementation is not too surprising, just adding the > required target methods. But there are some things worthy of mention: > > - amd-dbgapi does not support read watchpoints. If the core attempts > to insert a read (or access, which means read/write) watchpoint, > amd_dbgapi_target::insert_watchpoint returns an error. > > If we silently let the beneath target (linux-nat) install the read > watchpoint, it would be potentially confusing. Everything would look > fine to the user, but a read from the GPU would not be caught, so it > would look like the watchpoint doesn't work. > > There is a loophole though: read watchpoints created before the > runtime is loaded (and therefore the amd-dbgapi target is pushed) > will still be inserted. Only when execution stops, and the user > tries to resume again, will the check in > amd_dbgapi_target::insert_watchpoint be hit. > > Another option would be to allow the host read watchpoint to go > through, but warn that the reads from the AMD GPU device will not be > watched. We would need to be smart to avoid flooding the user with > warnings. But I decided to upstream the current ROCgdb behavior > first, we can always change it later. > > - When the amd-dbgapi target gets pushed, we create amd-dbgapi > watchpoints for any existing hardware write watchpoint location. > > - When the core asks the target to insert a watchpoint, we ask the > target beneath to insert it first. If the beneath target fails, we > return immediately with an error. > > - When the core asks to remove a watchpoint, we ask the target beneath > to to remove it first. Even if it fails, we still try to remove the > amd-dbgapi watchpoint. > > - When stopping after a watchpoint hit while the "precise-memory" > setting is not enabled, it is possible for the wave to stop a few > instructions later than the instruction that made the write that > triggered the watchpoint. We print a warning in that case, similar > to what we do when a memory violation happens while "precis-memory" > is disabled. > > Testing > ------- > > - Tests precise-memory-warning-watchpoint.exp and > watchpoint-at-end-of-shader.exp are more or less brought as-is from > downstream ROCgdb. I modified precise-memory-warning-watchpoint.exp > to watch a hipMalloc'ed region instead of a `__device__` global > variable. The latter doesn't work upstream, because we don't yet > support the DWARF constructs that describe the variable location. > > - I added test watchpoint-basic.exp with various simple cases to > exercises different code paths added by this patch. > > Differences from downstream ROCgdb > ---------------------------------- > > While extracting this code from ROCgdb, I made a few minor but possibly > significant (read: erroneous) changes. Those should be reviewed > carefully. I think that some code in ROCgdb was written at a time where > the amd-dbgapi target was always pushed at the very start of the > inferior execution, so assumptions were different. > > - The value type for the `amd_dbgapi_inferior_info::watchpoint_map` map > is now a structure, instead of an std::pair, just because it makes > the code more readable. > > - The insert_watchpoint and remove_watchpoint methods (and perhaps > others) now assume that if they are called, the runtime is in the > "enabled" state. > > - insert_initial_watchpoints has one more check (loc->owner->type != > bp_hardware_watchpoint), to filter out non-write watchpoints. > Otherwise, I think that we could mistakenly insert some write > watchpoints for some pre-existing read watchpoints. > > - Because it is possible for read watchpoints to be created before the > target is pushed, remove_watchpoint returns early if it sees that the > code asks for the removal of a read watchpoint, instead of asserting > "type == hw_write" (this was caught by the new test). > > - In ROCgdb, remove_watchpoint does: > > if (addr < it->first || (addr + len) > it->second.first) > return 1; > > I replaced it with some assertions. > > The first half of this condition should always be true, due to how > std::upper_bound works. > > For the second part: if the watchpoint was created successfully, it > is because it did fully cover the requested region (see > insert_one_watchpoint). I don't see why the core would ask us to > remove a watchpoint that wasn't successfully inserted. I am not 100% > sure about that one, there might be some edge cases where this is not > true. > I can't really see how this would be an issue, and your changes seem a reasonable assumption to me. I am happy going with this. > - I changed a manual free in stopped_by_watchpoint to a > gdb::unique_xmalloc_ptr, even though it changes nothing functionally. > > - I merged some conditions in amd_dbgapi_target_normal_stop. > Might be worth adding Laurent as co-auther, he did the initial implementation of this patch. With that, this looks good to me. I have tested it against the downstream testsuite as well, no issue reported there. Best, Lancelot. Approved-by: Lancelot Six > Change-Id: Ia15fb7434dc0c142a5a32997ada2e3a163c89f98 > --- > gdb/amd-dbgapi-target.c | 272 +++++++++++++++- > gdb/breakpoint.c | 4 +- > gdb/breakpoint.h | 4 + > .../precise-memory-warning-watchpoint.cpp | 43 +++ > .../precise-memory-warning-watchpoint.exp | 52 +++ > gdb/testsuite/gdb.rocm/rocm-test-utils.h | 12 + > .../gdb.rocm/watchpoint-at-end-of-shader.cpp | 44 +++ > .../gdb.rocm/watchpoint-at-end-of-shader.exp | 114 +++++++ > gdb/testsuite/gdb.rocm/watchpoint-basic.cpp | 69 ++++ > gdb/testsuite/gdb.rocm/watchpoint-basic.exp | 307 ++++++++++++++++++ > 10 files changed, 918 insertions(+), 3 deletions(-) > create mode 100644 gdb/testsuite/gdb.rocm/precise-memory-warning-watchpoint.cpp > create mode 100644 gdb/testsuite/gdb.rocm/precise-memory-warning-watchpoint.exp > create mode 100644 gdb/testsuite/gdb.rocm/watchpoint-at-end-of-shader.cpp > create mode 100644 gdb/testsuite/gdb.rocm/watchpoint-at-end-of-shader.exp > create mode 100644 gdb/testsuite/gdb.rocm/watchpoint-basic.cpp > create mode 100644 gdb/testsuite/gdb.rocm/watchpoint-basic.exp > > diff --git a/gdb/amd-dbgapi-target.c b/gdb/amd-dbgapi-target.c > index 471b7a7725ed..4e52683dc55a 100644 > --- a/gdb/amd-dbgapi-target.c > +++ b/gdb/amd-dbgapi-target.c > @@ -21,6 +21,7 @@ > #include "amd-dbgapi-target.h" > #include "amdgpu-tdep.h" > #include "async-event.h" > +#include "breakpoint.h" > #include "cli/cli-cmds.h" > #include "cli/cli-decode.h" > #include "cli/cli-style.h" > @@ -34,6 +35,8 @@ > #include "solib.h" > #include "target.h" > > +#include > + > /* When true, print debug messages relating to the amd-dbgapi target. */ > > static bool debug_amd_dbgapi = false; > @@ -227,6 +230,19 @@ struct amd_dbgapi_inferior_info > struct breakpoint *> > breakpoint_map; > > + /* Data associated to an inserted watchpoint. */ > + struct watchpoint_info > + { > + /* End address of the watched region. */ > + CORE_ADDR end_addr; > + > + /* ID returned by amd-dbgapi. */ > + amd_dbgapi_watchpoint_id_t id; > + }; > + > + /* Ordered map of inserted watchpoints. The key is the start address. */ > + std::map watchpoint_map; > + > /* List of pending events the amd-dbgapi target retrieved from the dbgapi. */ > std::list> wave_events; > > @@ -307,7 +323,12 @@ struct amd_dbgapi_target final : public target_ops > ULONGEST offset, ULONGEST len, > ULONGEST *xfered_len) override; > > + int insert_watchpoint (CORE_ADDR addr, int len, target_hw_bp_type type, > + expression *cond) override; > + int remove_watchpoint (CORE_ADDR addr, int len, target_hw_bp_type type, > + expression *cond) override; > bool stopped_by_watchpoint () override; > + std::vector stopped_data_addresses () override; > > bool stopped_by_sw_breakpoint () override; > bool stopped_by_hw_breakpoint () override; > @@ -711,13 +732,222 @@ amd_dbgapi_target::xfer_partial (enum target_object object, const char *annex, > return TARGET_XFER_OK; > } > > +/* Ask amd-dbgapi to insert a watchpoint in [ADDR, ADDR + len). > + > + Return 0 on success, 1 on failure. */ > + > +static int > +insert_one_watchpoint (amd_dbgapi_inferior_info *info, CORE_ADDR addr, int len) > +{ > + amd_dbgapi_watchpoint_id_t watch_id; > + > + if (amd_dbgapi_set_watchpoint (info->process_id, addr, len, > + AMD_DBGAPI_WATCHPOINT_KIND_STORE_AND_RMW, > + &watch_id) > + != AMD_DBGAPI_STATUS_SUCCESS) > + return 1; > + > + auto cleanup = make_scope_exit ([&] () > + { amd_dbgapi_remove_watchpoint (watch_id); }); > + > + /* A reduced range watchpoint may have been inserted, which would require > + additional watchpoints to be inserted to cover the requested range. > + > + For now, verify that the inserted watchpoint covers the requested range > + and error out if not. */ > + amd_dbgapi_global_address_t adjusted_address; > + > + if (amd_dbgapi_watchpoint_get_info (watch_id, > + AMD_DBGAPI_WATCHPOINT_INFO_ADDRESS, > + sizeof (adjusted_address), > + &adjusted_address) > + != AMD_DBGAPI_STATUS_SUCCESS > + || adjusted_address > addr) > + return 1; > + > + amd_dbgapi_size_t adjusted_size; > + > + if (amd_dbgapi_watchpoint_get_info (watch_id, > + AMD_DBGAPI_WATCHPOINT_INFO_SIZE, > + sizeof (adjusted_size), &adjusted_size) > + != AMD_DBGAPI_STATUS_SUCCESS > + || (adjusted_address + adjusted_size) < (addr + len)) > + return 1; > + > + using wp_info_t = amd_dbgapi_inferior_info::watchpoint_info; > + > + if (!(info->watchpoint_map.emplace (addr, wp_info_t {addr + len, watch_id}) > + .second)) > + return 1; > + > + cleanup.release (); > + return 0; > +} > + > +/* Insert watchpoints for all existing watchpoint locations associated to > + the program space of INFO. */ > + > +static void > +insert_initial_watchpoints (amd_dbgapi_inferior_info *info) > +{ > + gdb_assert (info->runtime_state == AMD_DBGAPI_RUNTIME_STATE_LOADED_SUCCESS); > + > + for (bp_location *loc : all_bp_locations ()) > + { > + /* Filter out other program spaces. */ > + if (loc->pspace != info->inf->pspace) > + continue; > + > + /* Filter out non-hardware watchpoints. */ > + if (loc->loc_type != bp_loc_hardware_watchpoint) > + continue; > + > + /* Filter out non-write watchpoints (access/read watchpoints might have > + been created before the runtime got loaded). */ > + if (loc->owner->type != bp_hardware_watchpoint) > + continue; > + > + if (insert_one_watchpoint (info, loc->address, loc->length) != 0) > + warning (_("Failed to insert existing watchpoint after loading " > + "runtime.")); > + } > +} > + > +int > +amd_dbgapi_target::insert_watchpoint (CORE_ADDR addr, int len, > + target_hw_bp_type type, expression *cond) > +{ > + amd_dbgapi_inferior_info &info > + = get_amd_dbgapi_inferior_info (current_inferior ()); > + > + /* The amd-dbgapi target is not pushed when the runtime is not active. */ > + gdb_assert (info.runtime_state == AMD_DBGAPI_RUNTIME_STATE_LOADED_SUCCESS); > + > + if (type != hw_write) > + { > + /* We only allow write watchpoints when GPU debugging is active. */ > + return 1; > + } > + > + if (int ret = beneath ()->insert_watchpoint (addr, len, type, cond); > + ret != 0) > + return ret; > + > + if (int ret = insert_one_watchpoint (&info, addr, len); > + ret != 0) > + { > + /* We failed to insert the GPU watchpoint, so remove the CPU watchpoint > + before returning an error. */ > + beneath ()->remove_watchpoint (addr, len, type, cond); > + return ret; > + } > + > + return 0; > +} > + > +int > +amd_dbgapi_target::remove_watchpoint (CORE_ADDR addr, int len, > + target_hw_bp_type type, > + expression *cond) > +{ > + amd_dbgapi_inferior_info &info > + = get_amd_dbgapi_inferior_info (current_inferior ()); > + > + /* The amd-dbgapi target is not pushed when the runtime is not active. */ > + gdb_assert (info.runtime_state == AMD_DBGAPI_RUNTIME_STATE_LOADED_SUCCESS); > + > + /* Try to remove the amd-dbgapi watchpoint even if the removal fails for the > + target beneath. */ > + int ret = beneath ()->remove_watchpoint (addr, len, type, cond); > + > + /* We don't allow non-write watchpoints (see the insert_watchpoints method) > + when the runtime is enabled (i.e. when the amd-dbgapi target is pushed). > + But there is a loophole: non-write watchpoints can still be created by the > + user before the runtime is enabled and the amd-dbgapi target is pushed. > + In that case, there won't be an amd-dbgapi watchpoint to remove, so just > + return. */ > + if (type != hw_write) > + return ret; > + > + /* Find the watchpoint id for the [addr, addr + len) range. */ > + auto it = info.watchpoint_map.upper_bound (addr); > + if (it == info.watchpoint_map.begin ()) > + return 1; > + > + std::advance (it, -1); > + > + /* Not a reference, so that we can reference wp_info after erasing *it. */ > + const auto [start_addr, wp_info] = *it; > + > + /* Since upper_bound finds the first element greater than ADDR, the previous > + element has to be less than or equal to ADDR. */ > + gdb_assert (start_addr <= addr); > + > + /* In insert_one_watchpoint, we ensured that the inserted watchpoint fully > + covered the requested range. It should be the same here. */ > + gdb_assert (addr + len <= wp_info.end_addr); > + > + info.watchpoint_map.erase (it); > + if (amd_dbgapi_remove_watchpoint (wp_info.id) != AMD_DBGAPI_STATUS_SUCCESS) > + return 1; > + > + return ret; > +} > + > bool > amd_dbgapi_target::stopped_by_watchpoint () > { > if (!ptid_is_gpu (inferior_ptid)) > return beneath ()->stopped_by_watchpoint (); > > - return false; > + amd_dbgapi_watchpoint_list_t watchpoints; > + if (amd_dbgapi_wave_get_info (get_amd_dbgapi_wave_id (inferior_ptid), > + AMD_DBGAPI_WAVE_INFO_WATCHPOINTS, > + sizeof (watchpoints), &watchpoints) > + != AMD_DBGAPI_STATUS_SUCCESS) > + return false; > + > + /* Ensure watchpoints.watchpoint_ids is freed on exit. */ > + gdb::unique_xmalloc_ptr > + watchpoint_ids_holder (watchpoints.watchpoint_ids); > + > + return watchpoints.count != 0; > +} > + > +std::vector > +amd_dbgapi_target::stopped_data_addresses () > +{ > + amd_dbgapi_inferior_info &info > + = get_amd_dbgapi_inferior_info (current_inferior ()); > + > + if (!ptid_is_gpu (inferior_ptid)) > + return beneath ()->stopped_data_addresses (); > + > + amd_dbgapi_watchpoint_list_t watchpoints = {}; > + if (amd_dbgapi_wave_get_info (get_amd_dbgapi_wave_id (inferior_ptid), > + AMD_DBGAPI_WAVE_INFO_WATCHPOINTS, > + sizeof (watchpoints), &watchpoints) > + != AMD_DBGAPI_STATUS_SUCCESS) > + return {}; > + > + /* Ensure watchpoints.watchpoint_ids is freed on exit. */ > + gdb::unique_xmalloc_ptr > + watchpoint_ids_holder (watchpoints.watchpoint_ids); > + > + std::vector watch_addr_hit; > + for (amd_dbgapi_watchpoint_id_t watch_id > + : gdb::make_array_view (watchpoints.watchpoint_ids, watchpoints.count)) > + { > + auto it = std::find_if (info.watchpoint_map.begin (), > + info.watchpoint_map.end (), > + [watch_id] (auto &wp) > + { return wp.second.id == watch_id; }); > + > + if (it != info.watchpoint_map.end ()) > + watch_addr_hit.push_back (it->first); > + } > + > + return watch_addr_hit; > } > > void > @@ -1031,9 +1261,12 @@ dbgapi_notifier_handler (int err, gdb_client_data client_data) > case AMD_DBGAPI_RUNTIME_STATE_LOADED_SUCCESS: > gdb_assert (info.runtime_state == AMD_DBGAPI_RUNTIME_STATE_UNLOADED); > info.runtime_state = runtime_state; > + > amd_dbgapi_debug_printf ("pushing amd-dbgapi target"); > info.inf->push_target (&the_amd_dbgapi_target); > > + insert_initial_watchpoints (&info); > + > /* The underlying target will already be async if we are running, but not if > we are attaching. */ > if (info.inf->process_target ()->is_async_p ()) > @@ -2297,6 +2530,42 @@ Warning: precise memory violation signal reporting is not enabled, reported\n\ > location may not be accurate. See \"show amdgpu precise-memory\".\n")); > } > > +/* Observer callback for normal_stop. Warn the user if a hardware watchpoint > + was hit but precise memory is not enabled. */ > + > +static void > +amd_dbgapi_target_normal_stop (bpstat *bs_list, int print_frame) > +{ > + if (bs_list == nullptr > + || !print_frame > + || !ptid_is_gpu (inferior_thread ()->ptid)) > + return; > + > + amd_dbgapi_inferior_info &info > + = get_amd_dbgapi_inferior_info (current_inferior ()); > + > + if (info.process_id == AMD_DBGAPI_PROCESS_NONE > + || info.precise_memory.enabled) > + return; > + > + bool found_hardware_watchpoint = false; > + > + for (bpstat *bs = bs_list; bs != nullptr; bs = bs->next) > + if (bs->breakpoint_at != nullptr > + && is_hardware_watchpoint (bs->breakpoint_at)) > + { > + found_hardware_watchpoint = true; > + break; > + } > + > + if (!found_hardware_watchpoint) > + return; > + > + gdb_printf (_("\ > +Warning: precise memory signal reporting is not enabled, watchpoint stop\n\ > +location may not be accurate. See \"show amdgpu precise-memory\".\n")); > +} > + > /* Style for some kinds of messages. */ > > static cli_style_option fatal_error_style > @@ -2529,6 +2798,7 @@ INIT_GDB_FILE (amd_dbgapi_target) > gdb::observers::inferior_exit.attach (amd_dbgapi_inferior_exited, "amd-dbgapi"); > gdb::observers::inferior_pre_detach.attach (amd_dbgapi_inferior_pre_detach, "amd-dbgapi"); > gdb::observers::thread_deleted.attach (amd_dbgapi_thread_deleted, "amd-dbgapi"); > + gdb::observers::normal_stop.attach (amd_dbgapi_target_normal_stop, "amd-dbgapi"); > > add_basic_prefix_cmd ("amdgpu", no_class, > _("Generic command for setting amdgpu flags."), > diff --git a/gdb/breakpoint.c b/gdb/breakpoint.c > index a4ccad32a8b9..da99ec27e19f 100644 > --- a/gdb/breakpoint.c > +++ b/gdb/breakpoint.c > @@ -2016,9 +2016,9 @@ is_breakpoint (const struct breakpoint *bpt) > || bpt->type == bp_dprintf); > } > > -/* Return true if BPT is of any hardware watchpoint kind. */ > +/* See breakpoint.h. */ > > -static bool > +bool > is_hardware_watchpoint (const struct breakpoint *bpt) > { > return (bpt->type == bp_hardware_watchpoint > diff --git a/gdb/breakpoint.h b/gdb/breakpoint.h > index 0d9111ba92e9..6b5dbcfe8a3c 100644 > --- a/gdb/breakpoint.h > +++ b/gdb/breakpoint.h > @@ -1085,6 +1085,10 @@ struct watchpoint : public breakpoint > > extern bool is_breakpoint (const struct breakpoint *bpt); > > +/* Return true if BPT is of any hardware watchpoint kind. */ > + > +extern bool is_hardware_watchpoint (const struct breakpoint *bpt); > + > /* Return true if BPT is of any watchpoint kind, hardware or > software. */ > > diff --git a/gdb/testsuite/gdb.rocm/precise-memory-warning-watchpoint.cpp b/gdb/testsuite/gdb.rocm/precise-memory-warning-watchpoint.cpp > new file mode 100644 > index 000000000000..dfa50be4b0da > --- /dev/null > +++ b/gdb/testsuite/gdb.rocm/precise-memory-warning-watchpoint.cpp > @@ -0,0 +1,43 @@ > +/* This testcase is part of GDB, the GNU debugger. > + > + Copyright 2021-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 . */ > + > +#include > + > +#include "rocm-test-utils.h" > + > +__global__ void > +kernel (int *ptr) > +{ > + for (int i = 0; i < 1000; ++i) > + (*ptr)++; > +} > + > +int *global_ptr; > + > +int > +main (int argc, char* argv[]) > +{ > + CHECK (hipMalloc (&global_ptr, sizeof (int))); > + CHECK (hipMemset (global_ptr, 0, sizeof (int))); > + > + /* Break here. */ > + hipLaunchKernelGGL (kernel, dim3 (1), dim3 (1), 0, 0, global_ptr); > + CHECK (hipDeviceSynchronize ()); > + > + CHECK (hipFree (global_ptr)); > + return 0; > +} > diff --git a/gdb/testsuite/gdb.rocm/precise-memory-warning-watchpoint.exp b/gdb/testsuite/gdb.rocm/precise-memory-warning-watchpoint.exp > new file mode 100644 > index 000000000000..9f8dac68047b > --- /dev/null > +++ b/gdb/testsuite/gdb.rocm/precise-memory-warning-watchpoint.exp > @@ -0,0 +1,52 @@ > +# Copyright 2021-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 when "amdgpu precise-memory" is off, hitting a watchpoint shows a > +# warning about the stop location maybe being inaccurate. > + > +load_lib rocm.exp > + > +require allow_hipcc_tests > + > +if { ![istarget "*-linux*"] } then { > + continue > +} > + > +standard_testfile .cpp > + > +if {[build_executable "failed to prepare" $testfile $srcfile {debug hip}]} { > + return > +} > + > +proc do_test { } { > + clean_restart > + gdb_load $::binfile > + > + with_rocm_gpu_lock { > + if {![runto kernel allow-pending qualified]} { > + fail "can't run to main" > + return > + } > + gdb_test "watch *((int *) global_ptr)" \ > + "Hardware watchpoint $::decimal: .*" \ > + "insert watchpoint" > + > + gdb_test "continue" \ > + "hit Hardware watchpoint $::decimal.*Warning: precise memory signal reporting is not enabled.*" \ > + "continue to watchpoint" > + } > +} > + > +do_test > diff --git a/gdb/testsuite/gdb.rocm/rocm-test-utils.h b/gdb/testsuite/gdb.rocm/rocm-test-utils.h > index 70d232b873b8..cbb1f6544617 100644 > --- a/gdb/testsuite/gdb.rocm/rocm-test-utils.h > +++ b/gdb/testsuite/gdb.rocm/rocm-test-utils.h > @@ -37,4 +37,16 @@ > } \ > while (0) > > +/* Ensure that all memory operations are completed before continuing, > + even when "precise-memory" is off. */ > + > +#define WAIT_MEM \ > + asm volatile (".if .amdgcn.gfx_generation_number < 10\n" \ > + " s_waitcnt 0\n" \ > + ".elseif .amdgcn.gfx_generation_number < 11\n" \ > + " s_waitcnt_vscnt null, 0\n" \ > + ".else\n" \ > + " s_wait_idle\n" \ > + ".endif") > + > #endif /* ROCM_TEST_UTILS_H */ > diff --git a/gdb/testsuite/gdb.rocm/watchpoint-at-end-of-shader.cpp b/gdb/testsuite/gdb.rocm/watchpoint-at-end-of-shader.cpp > new file mode 100644 > index 000000000000..e25c93e78489 > --- /dev/null > +++ b/gdb/testsuite/gdb.rocm/watchpoint-at-end-of-shader.cpp > @@ -0,0 +1,44 @@ > +/* This testcase is part of GDB, the GNU debugger. > + > + Copyright 2024-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 . */ > + > +#include > + > +#include "rocm-test-utils.h" > + > +__global__ void > +set_val (int *p, int v) > +{ > + *p = v; > +} > + > +int > +main () > +{ > + int *v; > + CHECK (hipMalloc (&v, sizeof (*v))); > + > + /* First dispatch to initialize the memory. */ > + set_val<<<1, 1>>> (v, 64); > + CHECK (hipDeviceSynchronize ()); > + > + /* Break here. */ > + set_val<<<1, 1>>> (v, 8); > + > + CHECK (hipDeviceSynchronize ()); > + > + return 0; > +} > diff --git a/gdb/testsuite/gdb.rocm/watchpoint-at-end-of-shader.exp b/gdb/testsuite/gdb.rocm/watchpoint-at-end-of-shader.exp > new file mode 100644 > index 000000000000..ff233c8b26a6 > --- /dev/null > +++ b/gdb/testsuite/gdb.rocm/watchpoint-at-end-of-shader.exp > @@ -0,0 +1,114 @@ > +# Copyright 2024-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 . > + > +# This test checks that a write watchpoint is reported to the debugger > +# when the memory is modified by the last statement of a shader. > + > +load_lib rocm.exp > + > +require allow_hipcc_tests > + > +standard_testfile .cpp > + > +if {[build_executable "failed to prepare" $testfile $srcfile {debug hip}]} { > + return > +} > + > +# This test is known to fail for some architectures when precise memory is > +# not enabled. This proc finds out if the current target is expected to have > +# a fail or not. > +proc target_has_xfail {} { > + # Check for targets where this test will fail without precise memory. > + set targets [find_amdgpu_devices] > + if { [llength $targets] == 0} { > + # Can't determine GPU type, don't set up xfail. The test will probably > + # not run correctly anyway. > + return 0 > + } > + > + # The test will run on GPU-0, so it should be the first of the list. > + set target [lindex $targets 0] > + > + # Extract the target family by removing the last 2 chars from the target > + # gfx number. > + set target_family [string range $target 0 end-2] > + verbose -log "Target family: $target_family" > + return [expr {[lsearch -exact {gfx10 gfx11 gfx12} $target_family] != -1}] > +} > + > +proc do_test {precise_memory has_xfail} { > + clean_restart > + gdb_load $::binfile > + > + with_rocm_gpu_lock { > + if {![runto [gdb_get_line_number "Break here."] allow-pending]} { > + return > + } > + > + gdb_test "p *v" "= 64" > + gdb_test "watch -l *v" "Hardware watchpoint $::decimal: -location \\*v" > + > + if {$precise_memory} { > + # For architectures that does not support precise memory, a warning > + # will be displayed. > + gdb_test "set amdgpu precise-memory $precise_memory" \ > + "(warning: AMDGPU precise memory access reporting could not be enabled\\.)?" > + } > + > + if { !$precise_memory && $has_xfail } { > + setup_xfail "*-*-*" > + } > + gdb_test "continue" \ > + [multi_line "Switching to thread $::decimal, lane 0.*" \ > + "" \ > + "Thread $::decimal \".*\" hit Hardware watchpoint $::decimal: -location \\*v" \ > + "" \ > + "Old value = 64" \ > + "New value = 8" \ > + ".*"] > + } > +} > + > +set has_xfail [target_has_xfail] > + > +# First check if we support precise-memory. > +set supports_precise_memory 0 > +clean_restart > +gdb_load $::binfile > +with_rocm_gpu_lock { > + if {![runto_main]} { > + return > + } > + > + gdb_test_multiple "set amdgpu precise-memory on" "" { > + -re -wrap "warning: AMDGPU precise memory access reporting could not be enabled\\." { > + set supports_precise_memory 0 > + pass $gdb_test_name > + } > + -re -wrap "^" { > + set supports_precise_memory 1 > + pass $gdb_test_name > + } > + } > +} > + > +foreach_with_prefix precise_memory {on off} { > + if { $precise_memory && !$supports_precise_memory } { > + unsupported "target does not support precise memory" > + continue > + } > + > + do_test $precise_memory $has_xfail > +} > diff --git a/gdb/testsuite/gdb.rocm/watchpoint-basic.cpp b/gdb/testsuite/gdb.rocm/watchpoint-basic.cpp > new file mode 100644 > index 000000000000..a89f33678f4b > --- /dev/null > +++ b/gdb/testsuite/gdb.rocm/watchpoint-basic.cpp > @@ -0,0 +1,69 @@ > +/* This testcase is part of GDB, the GNU debugger. > + > + 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 . */ > + > +#include > + > +#include "rocm-test-utils.h" > + > +__global__ void > +kernel (int *val1, int *val2) > +{ > + *val1 += 10; > + WAIT_MEM; > + *val2 += 100; > + WAIT_MEM; > + *val1 += 20; > + WAIT_MEM; > + *val2 += 200; > + WAIT_MEM; > + > + /* Some devices that don't support "precise memory" miss watchpoints when > + they would trigger near the end of the kernel. Execute a bunch of sleeps > + to make sure this doesn't happen. Just a handful of instructions should > + be enough, but this executes quickly anyway. */ > + for (int i = 0; i < 100000; ++i) > + __builtin_amdgcn_s_sleep (8); > +} > + > +/* Global pointers for the test to watch. */ > +int *global_ptr1; > +int *global_ptr2; > +int host_global = 5; > + > +int > +main () > +{ > + /* Break before runtime load. */ > + CHECK (hipMalloc (&global_ptr1, sizeof (int))); > + CHECK (hipMalloc (&global_ptr2, sizeof (int))); > + CHECK (hipMemset (global_ptr1, 0, sizeof (int))); > + CHECK (hipMemset (global_ptr2, 0, sizeof (int))); > + > + /* Break after malloc. */ > + kernel<<<1, 1>>> (global_ptr1, global_ptr2); > + CHECK (hipDeviceSynchronize ()); > + > + host_global += 12; > + > + /* Break before second launch. */ > + kernel<<<1, 1>>> (global_ptr1, global_ptr2); > + CHECK (hipDeviceSynchronize ()); > + > + CHECK (hipFree (global_ptr1)); > + CHECK (hipFree (global_ptr2)); > + return 0; > +} > diff --git a/gdb/testsuite/gdb.rocm/watchpoint-basic.exp b/gdb/testsuite/gdb.rocm/watchpoint-basic.exp > new file mode 100644 > index 000000000000..cd4b9c68f58a > --- /dev/null > +++ b/gdb/testsuite/gdb.rocm/watchpoint-basic.exp > @@ -0,0 +1,307 @@ > +# 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 . > + > +# Basic watchpoint tests for AMD GPU targets. > + > +load_lib rocm.exp > + > +require allow_hipcc_tests > + > +standard_testfile .cpp > + > +if {[build_executable "failed to prepare" $testfile $srcfile {debug hip}]} { > + return > +} > + > +proc continue_to_watchpoint_hit { old_value new_value test } { > + gdb_test "continue" \ > + [multi_line \ > + "hit Hardware watchpoint $::decimal:.*" \ > + "" \ > + "Old value = $old_value" \ > + "New value = $new_value" \ > + ".*"] \ > + $test > +} > + > +# Test inserting a watchpoint on a host variable before the runtime loads, and > +# hitting hit after the runtime loads. > + > +proc_with_prefix test_host_watchpoint_before_runtime_load {} { > + clean_restart $::testfile > + > + with_rocm_gpu_lock { > + if {![runto [gdb_get_line_number "Break before runtime load"]]} { > + return > + } > + > + gdb_test "watch host_global" \ > + "Hardware watchpoint $::decimal: .*" \ > + "set watchpoint on host_global" > + > + continue_to_watchpoint_hit 5 17 "continue to watchpoint hit" > + > + gdb_test "continue" \ > + "Inferior 1 .* exited normally.*" \ > + "continue to end" > + } > +} > + > +# Test inserting a watchpoint on a host variable after the runtime loads and > +# hitting hit. > + > +proc_with_prefix test_host_watchpoint_after_runtime_load {} { > + clean_restart $::testfile > + > + with_rocm_gpu_lock { > + if {![runto [gdb_get_line_number "Break after malloc"]]} { > + return > + } > + > + gdb_test "watch host_global" \ > + "Hardware watchpoint $::decimal: .*" \ > + "set watchpoint on host_global" > + > + continue_to_watchpoint_hit 5 17 "continue to watchpoint hit" > + > + gdb_test "continue" \ > + "Inferior 1 .* exited normally.*" \ > + "continue to end" > + } > +} > + > +# Test inserting a watchpoint before the kernel is launched, then hitting > +# it when the kernel runs. > + > +proc_with_prefix test_watchpoint_before_kernel {} { > + clean_restart $::testfile > + > + with_rocm_gpu_lock { > + if {![runto [gdb_get_line_number "Break after malloc"]]} { > + return > + } > + > + gdb_test "watch *((int *) global_ptr1)" \ > + "Hardware watchpoint $::decimal: .*" \ > + "set watchpoint on *ptr1" > + > + continue_to_watchpoint_hit 0 10 "continue to watchpoint hit 1" > + continue_to_watchpoint_hit 10 30 "continue to watchpoint hit 2" > + continue_to_watchpoint_hit 30 40 "continue to watchpoint hit 3" > + continue_to_watchpoint_hit 40 60 "continue to watchpoint hit 4" > + > + gdb_test "continue" \ > + "Inferior 1 .* exited normally.*" \ > + "continue to end" > + } > +} > + > +# Test inserting a watchpoint while stopped inside the kernel. > + > +proc_with_prefix test_watchpoint_inside_kernel {} { > + clean_restart $::testfile > + > + with_rocm_gpu_lock { > + if {![runto_main]} { > + return > + } > + > + gdb_breakpoint "kernel" allow-pending temporary > + gdb_test "continue" \ > + "hit Temporary breakpoint.*, kernel .*" \ > + "continue to kernel" > + > + gdb_test "watch *((int *) global_ptr2)" \ > + "Hardware watchpoint $::decimal: .*" \ > + "set watchpoint on *ptr2 from inside kernel" > + > + continue_to_watchpoint_hit 0 100 "continue to watchpoint hit 1" > + continue_to_watchpoint_hit 100 300 "continue to watchpoint hit 2" > + continue_to_watchpoint_hit 300 400 "continue to watchpoint hit 3" > + continue_to_watchpoint_hit 400 600 "continue to watchpoint hit 4" > + > + gdb_test "continue" \ > + "Inferior 1 .* exited normally.*" \ > + "continue to end" > + } > +} > + > +# Test removing a watchpoint while stopped inside the kernel, then > +# continuing to the end. > + > +proc_with_prefix test_remove_watchpoint_inside_kernel {} { > + clean_restart $::testfile > + > + with_rocm_gpu_lock { > + if {![runto [gdb_get_line_number "Break after malloc"]]} { > + return > + } > + > + gdb_test "watch *((int *) global_ptr1)" \ > + "Hardware watchpoint $::decimal: .*" \ > + "set watchpoint" > + > + continue_to_watchpoint_hit 0 10 "continue to watchpoint hit" > + gdb_test "with confirm off -- delete" "" "delete all breakpoints" > + > + gdb_test "continue" \ > + "Inferior 1 .* exited normally.*" \ > + "continue to end" > + } > +} > + > +# Test watchpoints on different memory locations. > + > +proc_with_prefix test_multiple_watchpoints {} { > + clean_restart $::testfile > + > + with_rocm_gpu_lock { > + if {![runto [gdb_get_line_number "Break after malloc"]]} { > + return > + } > + > + gdb_test "watch *((int *) global_ptr1)" \ > + "Hardware watchpoint $::decimal: .*" \ > + "set watchpoint on *ptr1" > + > + gdb_test "watch *((int *) global_ptr2)" \ > + "Hardware watchpoint $::decimal: .*" \ > + "set watchpoint on *ptr2" > + > + continue_to_watchpoint_hit 0 10 "continue to watchpoint hit 1" > + continue_to_watchpoint_hit 0 100 "continue to watchpoint hit 2" > + continue_to_watchpoint_hit 10 30 "continue to watchpoint hit 3" > + continue_to_watchpoint_hit 100 300 "continue to watchpoint hit 4" > + continue_to_watchpoint_hit 30 40 "continue to watchpoint hit 5" > + continue_to_watchpoint_hit 300 400 "continue to watchpoint hit 6" > + continue_to_watchpoint_hit 40 60 "continue to watchpoint hit 7" > + continue_to_watchpoint_hit 400 600 "continue to watchpoint hit 8" > + > + gdb_test "continue" \ > + "Inferior 1 .* exited normally.*" \ > + "continue to end" > + } > +} > + > +# Test disabling and enabling watchpoints. > + > +proc_with_prefix test_disable_enable_watchpoint {} { > + clean_restart $::testfile > + > + with_rocm_gpu_lock { > + if {![runto [gdb_get_line_number "Break after malloc"]]} { > + return > + } > + > + set wp1_num -1 > + set wp2_num -2 > + > + gdb_test_multiple "watch *((int *) global_ptr1)" "set watchpoint on *ptr1" { > + -re -wrap "Hardware watchpoint ($::decimal): .*" { > + set wp1_num $expect_out(1,string) > + pass $gdb_test_name > + } > + } > + > + gdb_test_multiple "watch *((int *) global_ptr2)" "set watchpoint on *ptr2" { > + -re -wrap "Hardware watchpoint ($::decimal): .*" { > + set wp2_num $expect_out(1,string) > + pass $gdb_test_name > + } > + } > + > + continue_to_watchpoint_hit 0 10 "continue to watchpoint hit 1" > + > + gdb_test_no_output "disable $wp1_num" "disable watchpoint" > + > + continue_to_watchpoint_hit 0 100 "continue to watchpoint hit 2" > + continue_to_watchpoint_hit 100 300 "continue to watchpoint hit 3" > + > + gdb_test_no_output "enable $wp1_num" "enable wp1" > + gdb_test_no_output "disable $wp2_num" "disable wp2" > + > + continue_to_watchpoint_hit 30 40 "continue to watchpoint hit 4" > + continue_to_watchpoint_hit 40 60 "continue to watchpoint hit 5" > + > + gdb_test_no_output "enable $wp2_num" "enable wp2" > + > + continue_to_watchpoint_hit 400 600 "continue to watchpoint hit 6" > + > + gdb_test "continue" \ > + "Inferior 1 .* exited normally.*" \ > + "continue to end" > + } > +} > + > +# Test that read and access watchpoints are rejected when GPU debugging is > +# active. The amd-dbgapi target only supports write watchpoints. WP_CMD > +# must either be "rwatch" or "awatch". > + > +proc_with_prefix test_non_write_watchpoint_rejected { wp_cmd } { > + clean_restart $::testfile > + > + with_rocm_gpu_lock { > + if {![runto [gdb_get_line_number "Break after malloc"]]} { > + return > + } > + > + gdb_test "$wp_cmd host_global" \ > + "Hardware .*watchpoint $::decimal: .*" \ > + "set watchpoint" > + > + gdb_test "continue" \ > + "Could not insert hardware watchpoint $::decimal.*Command aborted\\." \ > + "watchpoint rejected on resume" > + } > +} > + > +# Test setting a read/access watchpoint before the GPU runtime loads. The > +# watchpoint is set successfully on the CPU, but when the runtime loads, > +# we expect some behavior (error or warning) since the GPU only supports > +# write watchpoints. WP_CMD must either be "rwatch" or "awatch". > + > +proc_with_prefix test_non_write_watchpoint_before_runtime_load { wp_cmd } { > + clean_restart $::testfile > + > + with_rocm_gpu_lock { > + if {![runto [gdb_get_line_number "Break before runtime load"]]} { > + return > + } > + > + gdb_test "$wp_cmd host_global" \ > + "Hardware .*watchpoint $::decimal: .*" \ > + "set watchpoint" > + > + gdb_test "continue" "hit Hardware (read|access \\(read/write\\)) watchpoint.*Value = 5\r\n.*" \ > + "continue after setting watchpoint" > + > + gdb_test "continue" \ > + "Could not insert hardware watchpoint $::decimal.*Command aborted\\." \ > + "watchpoint rejected on resume" > + } > +} > + > +test_host_watchpoint_before_runtime_load > +test_host_watchpoint_after_runtime_load > +test_watchpoint_before_kernel > +test_watchpoint_inside_kernel > +test_remove_watchpoint_inside_kernel > +test_multiple_watchpoints > +test_disable_enable_watchpoint > + > +foreach_with_prefix wp_cmd { "rwatch" "awatch" } { > + test_non_write_watchpoint_rejected $wp_cmd > + test_non_write_watchpoint_before_runtime_load $wp_cmd > +} > -- > 2.52.0