#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"
#include "solib.h"
#include "target.h"
+#include <map>
+
/* When true, print debug messages relating to the amd-dbgapi target. */
static bool debug_amd_dbgapi = false;
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<CORE_ADDR, watchpoint_info> watchpoint_map;
+
/* List of pending events the amd-dbgapi target retrieved from the dbgapi. */
std::list<std::pair<ptid_t, target_waitstatus>> wave_events;
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<CORE_ADDR> stopped_data_addresses () override;
bool stopped_by_sw_breakpoint () override;
bool stopped_by_hw_breakpoint () override;
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<amd_dbgapi_watchpoint_id_t>
+ watchpoint_ids_holder (watchpoints.watchpoint_ids);
+
+ return watchpoints.count != 0;
+}
+
+std::vector<CORE_ADDR>
+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<amd_dbgapi_watchpoint_id_t>
+ watchpoint_ids_holder (watchpoints.watchpoint_ids);
+
+ std::vector<CORE_ADDR> 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
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 ())
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
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."),
|| 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
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. */
--- /dev/null
+/* 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 <http://www.gnu.org/licenses/>. */
+
+#include <hip/hip_runtime.h>
+
+#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;
+}
--- /dev/null
+# 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 <http://www.gnu.org/licenses/>.
+
+# 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
} \
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 */
--- /dev/null
+/* 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 <http://www.gnu.org/licenses/>. */
+
+#include <hip/hip_runtime.h>
+
+#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;
+}
--- /dev/null
+# 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 <http://www.gnu.org/licenses/>.
+
+# 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
+}
--- /dev/null
+/* 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 <http://www.gnu.org/licenses/>. */
+
+#include <hip/hip_runtime.h>
+
+#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;
+}
--- /dev/null
+# 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 <http://www.gnu.org/licenses/>.
+
+# 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
+}
projects / thirdparty / binutils-gdb.git / commitdiff
