From 5fc817ed78880516971296aa42ab49e05cd78139 Mon Sep 17 00:00:00 2001 From: Sarang Patrange Date: Wed, 8 Jul 2026 03:43:08 -0500 Subject: [PATCH] gdb/testsuite: add HIP graph-launch kernel debugging test Add gdb.rocm/hip-graph-launch test covering debugging of kernels dispatched via hipGraphLaunch (stream-captured graph, replayed). --- gdb/testsuite/gdb.rocm/hip-graph-launch.cpp | 106 ++++++++++++++++++++ gdb/testsuite/gdb.rocm/hip-graph-launch.exp | 94 +++++++++++++++++ 2 files changed, 200 insertions(+) create mode 100644 gdb/testsuite/gdb.rocm/hip-graph-launch.cpp create mode 100644 gdb/testsuite/gdb.rocm/hip-graph-launch.exp diff --git a/gdb/testsuite/gdb.rocm/hip-graph-launch.cpp b/gdb/testsuite/gdb.rocm/hip-graph-launch.cpp new file mode 100644 index 00000000000..569d7908dac --- /dev/null +++ b/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 . */ + +/* 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 +#include + +#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<<>> (result_ptr); + times_three<<>> (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; +} diff --git a/gdb/testsuite/gdb.rocm/hip-graph-launch.exp b/gdb/testsuite/gdb.rocm/hip-graph-launch.exp new file mode 100644 index 00000000000..c705478775f --- /dev/null +++ b/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 . + +# 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