gdb/testsuite: add HIP graph-launch kernel debugging test

gdb/testsuite/gdb.rocm/hip-graph-launch.cpp

@@ -0,0 +1,106 @@
+/* Copyright (C) 2026 Free Software Foundation, Inc.
+   Copyright (C) 2026 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/>.  */
+
+/* Build a HIP graph by capturing two kernel launches from a stream into a
+   single graph, instantiate it, and then launch (replay) the executable
+   graph several times with hipGraphLaunch.  This exercises the debugger's
+   ability to associate stops with the right kernel and source line when the
+   kernels reach the GPU through a graph launch -- the host submits a whole
+   pre-recorded graph of operations with hipGraphLaunch, which dispatches the
+   kernels -- rather than through a direct kernel<<<>>>() call.
+
+   The two kernels operate on the same buffer, in order, so each graph
+   replay computes out = (out + 1) * 3.  Starting from 0 this gives 3, 12
+   and 39 after the first, second and third replay respectively.  The
+   deterministic values let the .exp verify that each node executed in the
+   right order with the right data on every replay.  */
+
+#include <stdio.h>
+#include <hip/hip_runtime.h>
+
+#include "rocm-test-utils.h"
+
+/* Number of times the executable graph is replayed.  The .exp overrides
+   this via -DNUM_REPLAYS=... so the source and test stay in sync; the
+   default lets the program build on its own.  */
+#ifndef NUM_REPLAYS
+#define NUM_REPLAYS 3
+#endif
+
+/* First graph node.  */
+
+__global__ void
+add_one (int *out)
+{
+  int tid = threadIdx.x;
+  out[tid] = out[tid] + 1; /* break add_one */
+}
+
+/* Second graph node, run after add_one on the same buffer.  */
+
+__global__ void
+times_three (int *out)
+{
+  int tid = threadIdx.x;
+  out[tid] = out[tid] * 3; /* break times_three */
+}
+
+int
+main ()
+{
+  constexpr unsigned int num_elems = 1;
+
+  int *result_ptr;
+  CHECK (hipMalloc (&result_ptr, num_elems * sizeof (int)));
+  CHECK (hipMemset (result_ptr, 0, num_elems * sizeof (int)));
+
+  hipStream_t stream;
+  CHECK (hipStreamCreate (&stream));
+
+  /* Capture two kernel launches into a single graph.  */
+  CHECK (hipStreamBeginCapture (stream, hipStreamCaptureModeGlobal));
+  add_one<<<dim3 (1), dim3 (num_elems), 0, stream>>> (result_ptr);
+  times_three<<<dim3 (1), dim3 (num_elems), 0, stream>>> (result_ptr);
+  hipGraph_t graph;
+  CHECK (hipStreamEndCapture (stream, &graph));
+
+  /* Turn the graph into an executable graph.  */
+  hipGraphExec_t graph_exec;
+  CHECK (hipGraphInstantiate (&graph_exec, graph, nullptr, nullptr, 0));
+
+  /* This is the "launch graph" execution path: instead of the host
+     issuing each kernel<<<>>>(), hipGraphLaunch submits the whole graph and
+     re-dispatches both kernels on every replay.  The debugger should stop
+     in them on each launch.  */
+  for (int i = 0; i < NUM_REPLAYS; i++)
+    {
+      CHECK (hipGraphLaunch (graph_exec, stream));
+      CHECK (hipStreamSynchronize (stream));
+    }
+
+  int result;
+  CHECK (hipMemcpyDtoH (&result, result_ptr, sizeof (int)));
+  printf ("result is %d\n", result);
+
+  CHECK (hipGraphExecDestroy (graph_exec));
+  CHECK (hipGraphDestroy (graph));
+  CHECK (hipStreamDestroy (stream));
+  CHECK (hipFree (result_ptr));
+
+  return 0;
+}

gdb/testsuite/gdb.rocm/hip-graph-launch.exp

@@ -0,0 +1,94 @@
+# Copyright (C) 2026 Free Software Foundation, Inc.
+# Copyright (C) 2026 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 debugging kernels launched through a HIP graph (graph replay).
+#
+# The program captures two kernels (add_one then times_three) from a stream
+# into a single graph, instantiates it, and replays it NUM_REPLAYS times with
+# hipGraphLaunch.  Because the kernels reach the GPU through a graph launch
+# (the host submits a whole pre-recorded graph with hipGraphLaunch, which
+# dispatches the kernels) rather than a direct kernel<<<>>>() call, this
+# checks that the debugger still:
+#
+#   - resolves pending breakpoints in graph-launched kernels
+#   - associates each stop with the right kernel and source line
+#   - stops in each graph node, in order, on every replay (dispatch
+#     numbering / re-dispatch)
+#   - observes the correct per-replay data produced by each node
+
+load_lib rocm.exp
+
+require allow_hip_tests
+
+standard_testfile .cpp
+
+# Number of times the graph is replayed.  Passed to the C source via
+# -DNUM_REPLAYS so the source and test always stay in sync.
+set num_replays 3
+
+if {[build_executable "failed to prepare" $testfile $srcfile \
+	 [list debug hip additional_flags=-DNUM_REPLAYS=$num_replays]]} {
+    return
+}
+
+set add_one_line [gdb_get_line_number "break add_one"]
+set times_three_line [gdb_get_line_number "break times_three"]
+
+# Both kernel nodes are hit, in order, on every graph replay, and each
+# observes the value expected for that replay.
+proc_with_prefix test_replays_and_ordering {} {
+    with_rocm_gpu_lock {
+	clean_restart $::testfile
+
+	if {![runto_main]} {
+	    return
+	}
+
+	gdb_breakpoint "$::srcfile:$::add_one_line" -allow-pending
+	gdb_breakpoint "$::srcfile:$::times_three_line" -allow-pending
+
+	# out starts at 0.  Each replay computes out = (out + 1) * 3, so at the
+	# add_one stop out holds the value from the previous replay, and at the
+	# times_three stop it holds that value + 1.
+	set val 0
+	for {set i 1} {$i <= $::num_replays} {incr i} {
+	    with_test_prefix "replay $i" {
+		gdb_test "continue" \
+		    "Breakpoint .* add_one .*$::srcfile:$::add_one_line.*" \
+		    "stop in add_one"
+		gdb_test "print out\[0\]" " = $val" \
+		    "add_one sees value from previous replay"
+
+		gdb_test "continue" \
+		    "Breakpoint .* times_three .*$::srcfile:$::times_three_line.*" \
+		    "stop in times_three"
+		gdb_test "print out\[0\]" " = [expr {$val + 1}]" \
+		    "times_three sees add_one result"
+
+		set val [expr {($val + 1) * 3}]
+	    }
+	}
+
+	# After all replays the program prints the final result and exits.
+	gdb_test "continue" \
+	    "result is $val.*$::inferior_exited_re normally.*" \
+	    "continue to end, result $val"
+    }
+}
+
+test_replays_and_ordering