diff --git a/gdb/testsuite/gdb.rocm/runtime-core-attached.cpp b/gdb/testsuite/gdb.rocm/runtime-core-attached.cpp new file mode 100644 index 00000000000..d5201f60396 --- /dev/null +++ b/gdb/testsuite/gdb.rocm/runtime-core-attached.cpp @@ -0,0 +1,130 @@ +/* Copyright (C) 2023-2026 Free Software Foundation, Inc. + Copyright (C) 2023-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 . */ + +#include +#include +#include +#include +#include +#include +#include "gdb_watchdog.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) + +/* Pagefault kernel. In this testcase, OUT contains an address not reachable + by the GPU, triggering a page fault. */ + +__global__ void +pagefault_kernel (int *out) +{ + int local = 42; + *out = 8; +} + +/* This kernel will call abort (s_trap 2), which should cause the runtime to + generate a core dump. */ + +__global__ void +abort_kernel () +{ + int local = 42; + abort (); +} + +/* Secondary kernel, meant to run concurrently on a separate stream. This + kernel is meant to be running when the "main" kernel will generate an + exception. This is to ensure that GDB can load kernels which have raised an + exception (and entered the trap handler) and kernels which have not. */ + +__global__ void +aux_kernel () +{ + int local = 72; + + while (true) + __builtin_amdgcn_s_sleep (1); +} + +enum testcase_t +{ + memfault, + abort +}; + +int +main (int argc, char **argv) +{ + /* Make sure that the process terminates if the exception is not caught by + the ROCr runtime. */ + gdb_watchdog (30); + + if (argc != 2) + { + std::cerr + << "Usage: " << argv[0] << " pagefault|abort" << std::endl; + return EXIT_FAILURE; + } + + std::string teststr = argv[1]; + testcase_t test; + if (teststr == "pagefault") + test = testcase_t::memfault; + else if (teststr == "abort") + test = testcase_t::abort; + else + { + std::cerr << "Invalid test name \"" << teststr << "\"" << std::endl; + return EXIT_FAILURE; + } + + hipStream_t st1; + hipStream_t st2; + + CHECK (hipStreamCreate (&st1)); + CHECK (hipStreamCreate (&st2)); + + aux_kernel<<<1, 1, 0, st1>>> (); + + /* Make sure that the aux kernel gets time to start. */ + std::this_thread::sleep_for (std::chrono::seconds { 2 }); + + switch (test) + { + case testcase_t::memfault: + { + int *out = nullptr; + pagefault_kernel<<<1, 1, 0, st2>>> (out); + break; + } + case testcase_t::abort: + abort_kernel<<<1, 1, 0, st2>>> (); + break; + }; + + CHECK (hipDeviceSynchronize ()); +} diff --git a/gdb/testsuite/gdb.rocm/runtime-core-attached.exp b/gdb/testsuite/gdb.rocm/runtime-core-attached.exp new file mode 100644 index 00000000000..39f0b5a3f3b --- /dev/null +++ b/gdb/testsuite/gdb.rocm/runtime-core-attached.exp @@ -0,0 +1,118 @@ +# 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 core dump generation while GDB is attached. +# Verify that running a program under GDB with signal passthrough +# generates both host and GPU core dumps. + +load_lib rocm.exp + +require allow_rocm_core_tests +require allow_hip_tests + +standard_testfile .cpp + +if { [build_executable "failed to prepare" ${testfile} ${srcfile} \ + {debug hip}] } { + return -1 +} + +proc do_attached_test { fault } { + + set coredir [standard_output_file "coredir-$fault"] + remote_exec build "mkdir -p $coredir" + remote_exec build "rm -f $coredir/core*" + + clean_restart $::testfile + gdb_test "handle SIGABRT nostop pass" ".*" + gdb_test "handle SIGSEGV nostop pass" ".*" + gdb_test_no_output "set cwd $coredir" + gdb_test_no_output "set args $fault" + + gdb_run_cmd + + # Wait for program to crash and cores to be written. This can take + # a while for large core dumps. + with_timeout_factor 3 { + gdb_test_multiple "" "wait for program exit" { + -re "Program terminated with signal" { + exp_continue + } + -re "The program no longer exists" { + exp_continue + } + -re "$::gdb_prompt $" { + pass $gdb_test_name + } + timeout { + fail "$gdb_test_name (timeout)" + return + } + } + } + + set allcore_list [glob -nocomplain -directory $coredir "core*"] + set gpucore_list {} + set hostcore_list {} + + foreach core $allcore_list { + if {[string match "*.gpu" $core]} { + lappend gpucore_list $core + } else { + lappend hostcore_list $core + } + } + + if {[llength $gpucore_list] != 1} { + unsupported "single GPU core file must exist (found [llength $gpucore_list])" + return + } + + if {[llength $hostcore_list] != 1} { + unsupported "single host core file must exist (found [llength $hostcore_list])" + return + } + + set gpucore [lindex $gpucore_list 0] + set hostcore [lindex $hostcore_list 0] + + pass "generated both cores while gdb attached" + + set coremerge [gdb_find_coremerge] + set merged_core "$coredir/merged-$fault.core" + remote_exec build "$coremerge $merged_core $hostcore $gpucore" + + gdb_assert {[remote_file build exists $merged_core]} "merged cores" + + clean_restart $::testfile + gdb_test "target core $merged_core" \ + [multi_line \ + ".*" \ + "Core was generated by .*" \ + "Program terminated with signal SIG(ABRT|SEGV).*" \ + "Thread $::decimal \\(AMDGPU Wave \[^\r\n\]*\\).*"] \ + "load $fault merged corefile" + + gdb_test "info threads" "AMDGPU Wave.*" "verify GPU threads" +} + +with_rocm_gpu_lock { + foreach_with_prefix fault {pagefault abort} { + do_attached_test $fault + } +}