Add instruction stepping commands testcase

After a recent HIP change, the gdb.rocm/disassemble.exp testcase
started failing the "nexti" test:

 nexti
 0x00007ffff7f9b544 in __hip_get_thread_idx_x () at /opt/rocm-6.1.0-13445/include/hip/amd_detail/amd_hip_runtime.h:265
 265     __DEVICE__ unsigned int __hip_get_thread_idx_x() { return __ockl_get_local_id(0); }
 (gdb) FAIL: gdb.rocm/disassemble.exp: nexti

That now fails because:

1. - the testcase issues nexti when stopped at the "threadIdx.x" line,
     here:

       __global__ void kernel ()
       {
	 int tid = threadIdx.x;
	 tid += 1;
	 tid += 1;
       }

2. "threadIdx.x" now expands to a call to an inline function.

3. "nexti" doesn't skip inline functions.

4. The test is expecting GDB to stop at an instruction in the middle
of a line (the current frame output starts with an hex number for
current the PC address).

The current "nexti" test is, I argue, pretty useless, it only
basically checks that the command doesn't crash.

This commit replaces that test with a new testcase, that exercises
both "stepi" and "nexti".  The main difference between these two
commands is that nexti steps over function calls, while stepi does
not.  So the testcase issues multiple stepi/nexti commands until one
of these happen:

 - stepi - until we've entered a called function.

 - nexti - until we've stepped over the called function and reached a
   different line.

Note, the generic part of the testsuite already has tests for
something similar to this, in gdb.base/step-test.exp, that we should
be able to exercise with --target_board=hip.

gdb.rocm/disassemble.exp now passes, with the "bad test" removed.

Change-Id: Iaefa2629a2ac030913c3779add56ab2bf9b1f6bc

gdb/testsuite/gdb.rocm/disassemble.exp

@@ -93,12 +93,6 @@ gdb_test_sequence "x/3i kernel" "disassemble 3 instructions" [list \
     "\\s+$hex\\s+<kernel\\(\\)\\+?($decimal)?>:(\\s*)"
 ]
 
-# Check "nexti".
-# Sample vega20 output:
-#
-#  0x00007fa50e601664      27          int tid = threadIdx.x;
-gdb_test "nexti" "$hex\\s+$decimal.+"
-
 # Check set disassemble-next-line off followed by n.
 # Sample vega20 output:
 # There should not be any disassemble code printed identified by the absence of the character "=>"

gdb/testsuite/gdb.rocm/instruction-stepping-commands.cpp

@@ -0,0 +1,58 @@
+/* Copyright (C) 2024 Free Software Foundation, Inc.
+   Copyright (C) 2024 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 <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)
+
+/* Make sure the function isn't inlined, for the nexti test.  */
+__device__ static __attribute__ ((noinline))
+int
+return_zero ()
+{
+  return 0;
+}
+
+__global__ void
+kernel ()
+{
+  int var = return_zero ();
+  var += 1; /* next line */
+  var += 1;
+}
+
+int
+main (int argc, char* argv[])
+{
+  kernel<<<1, 1>>> ();
+  CHECK (hipDeviceSynchronize ());
+  return 0;
+}

gdb/testsuite/gdb.rocm/instruction-stepping-commands.exp

@@ -0,0 +1,84 @@
+# Copyright (C) 2024 Free Software Foundation, Inc.
+# Copyright (C) 2024 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/>.
+
+# Basic test for the instruction stepping commands: stepi/nexti.
+
+load_lib rocm.exp
+
+require allow_hipcc_tests
+
+standard_testfile .cpp
+
+if {[prepare_for_testing "failed to prepare ${testfile}" $testfile $srcfile {debug hip}]} {
+    return -1
+}
+
+# Issue CMD repeatedly until UNTIL_REGEXP matches the GDB output.  CMD
+# can be either "stepi" or "nexti".
+proc test_step_until {cmd until_regexp} {
+
+    if ![runto kernel allow-pending qualified message] {
+	return
+    }
+
+    # Number of instructions we've single-stepped.
+    set insn_count 0
+
+    # A reasonable limit, so that we don't loop forever if something
+    # goes wrong.
+    set insn_count_max 100
+
+    gdb_test_multiple $cmd "" {
+	-re -wrap "$::hex\\s+$::decimal\\s+int var = return_zero \\(\\);" {
+	    # Note: we expect to see a leading PC address as we're
+	    # stopping at instructions that map to the middle of a
+	    # source line.
+
+	    incr insn_count
+	    verbose -log "insn_count = $insn_count"
+
+	    if {$insn_count >= $insn_count_max} {
+		fail "$gdb_test_name (too many steps)"
+	    } else {
+		send_gdb "$cmd\n"
+		exp_continue
+	    }
+	}
+	-re -wrap "$until_regexp" {
+	    # We should have seen at least one instruction being
+	    # single-stepped.
+	    gdb_assert {$insn_count > 0} $gdb_test_name
+	}
+    }
+}
+
+with_rocm_gpu_lock {
+
+# With stepi, stepping eventually stops at the start of the function
+# called by the kernel entry point.
+with_test_prefix "stepi" {
+    test_step_until "stepi" "^return_zero.*"
+}
+
+# With nexti, stepping steps over function calls, and eventually stops
+# at the next line.
+with_test_prefix "nexti" {
+    test_step_until "nexti" "^$::decimal\\s+.*next line.*"
+}
+
+}