[PATCH 7/8] Add tests for s_endpgm handling
Pedro Alves
pedro@palves.net
Thu Dec 14 20:22:37 GMT 2023
[Not for commit. This won't work with current upstream, unfortunately.]
Check that a wave can halt at an s_endpgm instruction by
single-stepping or displaced stepping the instruction preceding the
s_endpgm.
Check that a wave can single-step s_endpgm, and that it can step over
a breakpoint placed at an s_endpgm by displaced-stepping the
instruction. Test all three cases:
- no step-over (stepping without a breakpoint installed)
- in-line step-over
- displaced step-over
Check the same with "set scheduler-locking on".
Check that GDB always prints the exited wave's ID when aborting a
command due to thread exit, and that it prints a valid ID with no "?"
in it.
This is named gdb.rocm/step-over-kernel-exit.cpp and not
gdb.rocm/s_endpgm.cpp because we will most probably want to extend
this to test s_sendmsg deallow vgprs before s_endpgm as well.
Co-Authored-By: Laurent Morichetti <laurent.morichetti@amd.com>
Co-Authored-By: Simon Marchi <simon.marchi@efficios.com>
Change-Id: I6db617ac009383698e1c66744d68e70b1d1ca90f
---
.../gdb.rocm/step-over-kernel-exit.cpp | 48 ++++++++
.../gdb.rocm/step-over-kernel-exit.exp | 108 ++++++++++++++++++
2 files changed, 156 insertions(+)
create mode 100644 gdb/testsuite/gdb.rocm/step-over-kernel-exit.cpp
create mode 100644 gdb/testsuite/gdb.rocm/step-over-kernel-exit.exp
diff --git a/gdb/testsuite/gdb.rocm/step-over-kernel-exit.cpp b/gdb/testsuite/gdb.rocm/step-over-kernel-exit.cpp
new file mode 100644
index 00000000000..61f1b431df1
--- /dev/null
+++ b/gdb/testsuite/gdb.rocm/step-over-kernel-exit.cpp
@@ -0,0 +1,48 @@
+/* Copyright (C) 2023 Free Software Foundation, Inc.
+ Copyright (C) 2023 Advanced Micro Devices, Inc. All rights reserved.
+
+ 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 <http://www.gnu.org/licenses/>.
+*/
+
+#include <stdio.h>
+#include <hip/hip_runtime.h>
+
+#define CHECK(cmd) \
+ do \
+ { \
+ hipError_t error = cmd; \
+ if (error != hipSuccess) \
+ { \
+ fprintf (stderr, "error: '%s'(%d) at %s:%d\n", \
+ hipGetErrorString (error), error, \
+ __FILE__, __LINE__); \
+ exit (EXIT_FAILURE); \
+ } \
+ } while (0)
+
+__global__ void
+kernel ()
+{
+ asm ("before_s_endpgm_insn: s_nop 0");
+ asm ("s_endpgm_insn: s_endpgm"); /* set breakpoint here */
+}
+
+int
+main (int argc, char **argv)
+{
+ kernel<<<1, 1>>> ();
+ CHECK (hipDeviceSynchronize ());
+}
diff --git a/gdb/testsuite/gdb.rocm/step-over-kernel-exit.exp b/gdb/testsuite/gdb.rocm/step-over-kernel-exit.exp
new file mode 100644
index 00000000000..484298ffa3e
--- /dev/null
+++ b/gdb/testsuite/gdb.rocm/step-over-kernel-exit.exp
@@ -0,0 +1,108 @@
+# Copyright (C) 2023 Free Software Foundation, Inc.
+# Copyright (C) 2023 Advanced Micro Devices, Inc. All rights reserved.
+
+# 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 <http://www.gnu.org/licenses/>.
+
+# Test stopping at and single-stepping and displaced-stepping an
+# s_endpgm instruction.
+
+load_lib rocm.exp
+
+standard_testfile .cpp
+
+require allow_hipcc_tests
+
+if { [build_executable "failed to prepare" \
+ $testfile $srcfile {debug hip}] == -1 } {
+ return -1
+}
+
+# Test stepping over an s_endpgm instruction.
+#
+# STEP_OVER_MODE can be one of:
+#
+# - none: don't put a breakpoint on the s_endpgm instruction.
+
+# - inline: put a breakpoint on the s_endpgm instruction, and use
+# in-line stepping to step over it (disable displaced-stepping).
+
+# - displaced: same, but use displaced stepping.
+#
+# SCHEDLOCK can be "on" or "off".
+
+proc do_test { step_over_mode schedlock } {
+ with_rocm_gpu_lock {
+ clean_restart $::binfile
+
+ if { $step_over_mode == "none" } {
+ # Nothing to do.
+ } elseif { $step_over_mode == "inline" } {
+ gdb_test_no_output "set displaced-stepping off"
+ } elseif { $step_over_mode == "displaced" } {
+ gdb_test_no_output "set displaced-stepping on"
+ } else {
+ error "Invalid step_over_mode value: $step_over_mode"
+ }
+
+ if ![runto_main] {
+ fail "can't run to main"
+ return -1
+ }
+
+ # Put a breakpoint on the instruction before s_endpgm,
+ # continue to it.
+ gdb_breakpoint "before_s_endpgm_insn" allow-pending
+ gdb_continue_to_breakpoint "before_s_endpgm_insn"
+
+ gdb_test_no_output "set scheduler-locking $schedlock"
+
+ gdb_test "stepi" \
+ "\"s_endpgm_insn: s_endpgm\".*" \
+ "single-step instruction before s_endpgm"
+
+ gdb_test "x/i \$pc" \
+ "$::hex <\[^\r\n\]*>:\[ \t\]+s_endpgm.*" \
+ "stopped at s_endpgm"
+
+ # If testing a step-over is requested, place a breakpoint at
+ # the current instruction to force a step-over.
+ if { $step_over_mode != "none" } {
+ gdb_test "break s_endpgm_insn" "Breakpoint $::decimal at $::hex.*"
+ }
+
+ set d $::decimal
+ set wave_target_id_re "AMDGPU Wave $d:$d:$d:1 \\(0,0,0\\)/0"
+
+ set selected_thread_before \
+ [get_integer_valueof "\$_thread" 0 "get selected thread before"]
+
+ gdb_test "stepi" \
+ "\r\n\[$wave_target_id_re exited\]\r\nCommand aborted, thread exited\\." \
+ "single-step s_endpgm"
+
+ # Check that the selected thread didn't change, and that GDB
+ # manages to print the exited wave's target ID properly.
+ gdb_test "thread" \
+ "\r\n\[Current thread is $selected_thread_before \\($wave_target_id_re\\) \\(exited\\)\]" \
+ "exited wave target id"
+ }
+}
+
+foreach_with_prefix step_over_mode {none inline displaced} {
+ foreach_with_prefix schedlock {off on} {
+ do_test $step_over_mode $schedlock
+ }
+}
--
2.43.0
More information about the Gdb-patches
mailing list