[PATCH v4 37/44] testsuite, sycl: add test for backtracing inside a kernel

[Markus Metzger <markus.t.metzger@intel.com>]

From: Tankut Baris Aktemur <tankut.baris.aktemur@intel.com>

Add SYCL test for checking the call stack inside a kernel, including
inlined functions.

Co-authored-by: Natalia Saiapova <natalia.saiapova@intel.com>
---
 gdb/testsuite/gdb.sycl/call-stack.cpp |  92 ++++++++++++++
 gdb/testsuite/gdb.sycl/call-stack.exp | 176 ++++++++++++++++++++++++++
 2 files changed, 268 insertions(+)
 create mode 100644 gdb/testsuite/gdb.sycl/call-stack.cpp
 create mode 100644 gdb/testsuite/gdb.sycl/call-stack.exp

diff --git a/gdb/testsuite/gdb.sycl/call-stack.cpp b/gdb/testsuite/gdb.sycl/call-stack.cpp
new file mode 100644
index 00000000000..516dd3e3324
--- /dev/null
+++ b/gdb/testsuite/gdb.sycl/call-stack.cpp
@@ -0,0 +1,92 @@
+/* This testcase is part of GDB, the GNU debugger.
+
+   Copyright 2019-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 <sycl/sycl.hpp>
+#include <iostream>
+#include "sycl-util.cpp"
+
+int
+fourth (int x4, int y4)
+{
+  return x4 * y4; /* ordinary-fourth-loc */
+}
+
+int
+third (int x3, int y3)
+{
+  return fourth (x3 + 5, y3 * 3) + 30; /* ordinary-third-loc */
+}
+
+int
+second (int x2, int y2)
+{
+  return third (x2 + 5, y2 * 3) + 30; /* ordinary-second-loc */
+}
+
+int
+first (int x1, int y1)
+{
+  int result = second (x1 + 5, y1 * 3); /* ordinary-first-loc */
+  return result + 30; /* kernel-function-return */
+}
+
+__attribute__((always_inline))
+int
+inlined_second (int x, int y)
+{
+  return x * y; /* inlined-inner-loc */
+}
+
+__attribute__((always_inline))
+int
+inlined_first (int num1, int num2)
+{
+  int result = inlined_second (num1 + 5, num2 * 3); /* inlined-middle-loc */
+  return result + 30;
+}
+
+int
+main (int argc, char *argv[])
+{
+  int data[3] = {7, 8, 9};
+
+  { /* Extra scope enforces waiting on the kernel.  */
+    sycl::queue deviceQueue {get_sycl_queue (argc, argv)};
+    sycl::buffer<int, 1> buf {data, sycl::range<1> {3}};
+
+    deviceQueue.submit ([&] (sycl::handler& cgh)
+      {
+	auto numbers = buf.get_access<sycl::access::mode::read_write> (cgh);
+
+	cgh.single_task ([=] ()
+	  {
+	    int ten = numbers[1] + 2;
+	    int four = numbers[2] - 5;
+	    int fourteen = ten + four;
+	    numbers[0] = first (fourteen + 1, 3); /* ordinary-outer-loc */
+	    numbers[1] = inlined_first (10, 2); /* inlined-outer-loc */
+	    numbers[2] = first (3, 4); /* another-call */
+	  });
+      });
+  }
+
+  std::cout << "Result is " << data[0] << " "
+	    << data[1] << " " << data[2] << std::endl;
+  /* Expected: 210 120 126 */
+
+  return 0; /* end-of-program */
+}
diff --git a/gdb/testsuite/gdb.sycl/call-stack.exp b/gdb/testsuite/gdb.sycl/call-stack.exp
new file mode 100644
index 00000000000..73e59dfeab7
--- /dev/null
+++ b/gdb/testsuite/gdb.sycl/call-stack.exp
@@ -0,0 +1,176 @@
+# Copyright 2019-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/>.
+#
+# Tests GDBs support for SYCL when there are function calls inside
+# the kernel.
+
+load_lib sycl.exp
+
+standard_testfile .cpp
+
+set sycl_device_list [init_sycl_devices_list]
+if {[llength $sycl_device_list] == 0} {
+    unsupported "target does not support SYCL"
+    return
+}
+
+if {[build_executable "failed to compile $srcfile" "$binfile" "$srcfile" \
+	 {sycl debug}]} {
+    return
+}
+
+# Return the current line number.
+proc get_current_line {} {
+    global decimal
+    gdb_test_multiple "info line" "get current line" {
+	-re -wrap "Line ($decimal).*" {
+	    pass $gdb_test_name
+	    return $expect_out(1,string)
+	}
+	-re -wrap "" {
+	    fail $gdb_test_name
+	    return 0
+	}
+    }
+}
+
+proc test_call_stack {device} {
+    global srcfile valnum_re decimal inferior_exited_re
+
+    set fourth_loc [gdb_get_line_number "ordinary-fourth-loc"]
+    set third_loc [gdb_get_line_number "ordinary-third-loc"]
+    set second_loc [gdb_get_line_number "ordinary-second-loc"]
+    set first_loc [gdb_get_line_number "ordinary-first-loc"]
+    set outer_loc [gdb_get_line_number "ordinary-outer-loc"]
+    set inlined_inner_loc [gdb_get_line_number "inlined-inner-loc"]
+    set inlined_middle_loc [gdb_get_line_number "inlined-middle-loc"]
+    set inlined_outer_loc [gdb_get_line_number "inlined-outer-loc"]
+
+    set fill "\[^\r\n\]*"
+
+    set fourth_desc "fourth \\(x4=$fill, y4=$fill\\) at ${fill}$srcfile:$fourth_loc"
+    set third_desc "third \\(x3=$fill, y3=$fill\\) at ${fill}$srcfile:$third_loc"
+    set second_desc "second \\(x2=$fill, y2=$fill\\) at ${fill}$srcfile:$second_loc"
+    set first_desc "first \\(x1=$fill, y1=$fill\\) at ${fill}$srcfile:$first_loc"
+    set outer_desc "${fill}operator\\(\\)${fill} at ${fill}$srcfile:$outer_loc"
+    set inlined_inner_desc \
+	"inlined_second ${fill} at ${fill}$srcfile:$inlined_inner_loc"
+    set inlined_middle_desc \
+	"inlined_first ${fill} at ${fill}$srcfile:$inlined_middle_loc"
+    set inlined_outer_desc \
+	"${fill}operator\\(\\)${fill} at ${fill}$srcfile:$inlined_outer_loc"
+
+    gdb_breakpoint "$srcfile:$fourth_loc"
+    gdb_continue_to_breakpoint "innermost-body" ".*$srcfile:$fourth_loc.*"
+
+    # Limit the backtrace to 5 frames because frame #5
+    # and beyond are implementation-specific to the SYCL runtime.
+    gdb_test "backtrace 5" [multi_line \
+				"#0${fill} $fourth_desc" \
+				"#1${fill} $third_desc" \
+				"#2${fill} $second_desc" \
+				"#3${fill} $first_desc" \
+				"#4${fill} $outer_desc.*"] \
+	"first backtrace"
+
+    # Test inlined function calls.
+    gdb_breakpoint $inlined_inner_loc
+
+    gdb_continue_to_breakpoint "inlined-body" ".*$srcfile:$inlined_inner_loc.*"
+
+    gdb_test "backtrace 3" [multi_line \
+				"#0${fill} $inlined_inner_desc" \
+				"#1${fill} $inlined_middle_desc" \
+				"#2${fill} $inlined_outer_desc.*"] \
+	"backtrace for inlined calls"
+
+    delete_breakpoints
+
+    # Now we will stop at the beginning of prologue of the fourth function
+    # and instruction step through the function until it returns back
+    # to the third.
+    gdb_breakpoint "*fourth"
+
+    gdb_test "continue" "fourth.*$srcfile.*" "continue to fourth prologue"
+    set i 0
+    set current_line [get_current_line]
+    set fourth_prologue_line $current_line
+
+    # Update description to not include arguments.
+    set third_desc "third ${fill} at ${fill}$srcfile:$third_loc"
+    set second_desc "second ${fill} at ${fill}$srcfile:$second_loc"
+    set first_desc "first ${fill} at ${fill}$srcfile:$first_loc"
+
+    # Print the current instruction and framedesc for logging purposes.
+    gdb_test "display/i \$pc"
+    gdb_test "display/x \$framedesc"
+
+    # Check the backtrace at each instruction until the return.  We do not
+    # check the args here, as they might be invalid at prologue and epilogue.
+    # Also check that there are no additional messages after
+    # the backtrace except "(More stack frames follow...)".
+    #
+    # Restrict the number of iterations to avoid an infinite loop in case
+    # of problems.
+    while {($current_line == $fourth_prologue_line
+	    || $current_line == $fourth_loc)
+	   && $i < 100} {
+	with_test_prefix "iteration $i" {
+	    if {[require_sycl_device "$device" "gpu" "Intel*"]} {
+		set fourth_desc "fourth ${fill} at ${fill}$srcfile:$current_line"
+		gdb_test "backtrace 6" [multi_line \
+					    "#0${fill} $fourth_desc" \
+					    "#1${fill} $third_desc" \
+					    "#2${fill} $second_desc" \
+					    "#3${fill} $first_desc" \
+					    "#4${fill} main${fill}operator${fill}lambda${fill} at .*" \
+					    "#5${fill}(\r\n\\(More stack frames follow\\.\\.\\.\\))?"] \
+		    "backtrace in fourth"
+	    } else {
+		# On CPU the backtrace in prologue might include
+		# additional RT specific frames.  Do not assume any
+		# frame numbers and do a deeper backtrace.  Do not
+		# expect the line number at fourth.  We expect to see
+		# our frames somewhere in the middle.
+		set fourth_desc "fourth ${fill} at ${fill}$srcfile:$decimal"
+		gdb_test "backtrace 10" [multi_line \
+					     "${fill} $fourth_desc" \
+					     "${fill} $third_desc" \
+					     "${fill} $second_desc" \
+					     "${fill} $first_desc" \
+					     "${fill} main${fill}operator${fill}lambda${fill} at .*"] \
+		    "backtrace in fourth"
+	    }
+	    gdb_test "with scheduler-locking on -- stepi"
+	    incr i
+	    set current_line [get_current_line]
+	}
+    }
+
+    # Disable printing of PC and FRAMEDESC.
+    gdb_test "undisplay 1-2"
+}
+
+foreach device $sycl_device_list {
+    sycl_with_intelgt_lock $device {
+	clean_restart $testfile
+
+	if {![sycl_start $device]} {
+	    continue
+	}
+
+	test_call_stack "$device"
+    }
+}
-- 
2.43.0

________________________________________
Intel Deutschland GmbH 

Registered Address: Dornacher Strasse 1, 85622 Feldkirchen, Germany 

Tel: +49 (89) 99143-0 

www.intel.de 

Managing Directors: Candice Moore, Jeffrey Schneiderman, Ramachandran Sitaraman

Chairperson of the Supervisory Board: Sonja Pierer

Registered Seat: Munich Commercial Register B: Amtsgericht Munich HRB 186928

This e-mail and any attachments may contain confidential material for
the sole use of the intended recipient(s). Any review or distribution
by others is strictly prohibited. If you are not the intended
recipient, please contact the sender and delete all copies.