[PATCH 0/6] Fix generating core files when GPU waves are present

[PATCH 0/6] Fix generating core files when GPU waves are present

[Simon Marchi]

The thing that sparked this series is that we identified a problem in
rocgdb with the recently introduced get_ptid_for_slash_proc function.
When GPU waves are present, and especially when a GPU wave is selected,
get_ptid_for_slash_proc will return the ptid of such a wave to be used
for /proc accesses.  All GPU waves have `lwp == 1`, which means GDB
attempts to read from /proc/1, which does not end well.

While looking for a way to expose this problem with what we currently
have upstream, I found that trying to generate a core dump while GPU
waves exist fails quite badly.  The get_ptid_for_slash_proc problem is a
reason, but there are more issues.

The goal of this series is to make it possible to generate a host core
dump while GPU waves are active, and possibly one of those is selected.
It is not a goal of this series to dump the state of GPU waves.
Downstream rocgdb has this functionality, and we want to send it
upstream eventually, just not in this series.

Shahab Vahedi (1):
  gdb/amd-dbgapi: defer TARGET_OBJECT_AUXV xfers to the beneath target

Simon Marchi (5):
  gdb/testsuite: make with_rocm_gpu_lock transparent to return
  gdb/corefile: handle no signalled thread when writing NT_GDB_TDESC
  gdb/amd-dbgapi: handle regno == -1 in fetch_registers and
    store_registers
  gdb/linux-tdep: don't write core file notes for non-Linux threads
  gdb/linux-tdep: ignore non-Linux threads in get_ptid_for_slash_proc

 gdb/amd-dbgapi-target.c                       | 109 +++++++++------
 gdb/elf-none-tdep.c                           |   7 +-
 gdb/gcore.c                                   |  27 ++--
 gdb/gcore.h                                   |   9 +-
 gdb/linux-tdep.c                              |  64 ++++++---
 gdb/testsuite/gdb.rocm/auxv-wave-selected.cpp |  42 ++++++
 gdb/testsuite/gdb.rocm/auxv-wave-selected.exp |  65 +++++++++
 gdb/testsuite/gdb.rocm/gcore-host.cpp         |  42 ++++++
 gdb/testsuite/gdb.rocm/gcore-host.exp         | 125 ++++++++++++++++++
 .../gdb.testsuite/with-rocm-gpu-lock.exp      |  90 +++++++++++++
 .../gcore-exited-selected-thread.c            |  51 +++++++
 .../gcore-exited-selected-thread.exp          |  83 ++++++++++++
 gdb/testsuite/lib/rocm.exp                    |  20 ++-
 13 files changed, 661 insertions(+), 73 deletions(-)
 create mode 100644 gdb/testsuite/gdb.rocm/auxv-wave-selected.cpp
 create mode 100644 gdb/testsuite/gdb.rocm/auxv-wave-selected.exp
 create mode 100644 gdb/testsuite/gdb.rocm/gcore-host.cpp
 create mode 100644 gdb/testsuite/gdb.rocm/gcore-host.exp
 create mode 100644 gdb/testsuite/gdb.testsuite/with-rocm-gpu-lock.exp
 create mode 100644 gdb/testsuite/gdb.threads/gcore-exited-selected-thread.c
 create mode 100644 gdb/testsuite/gdb.threads/gcore-exited-selected-thread.exp


base-commit: 835aee81fa93e587eb9ac73a79e911361ec46d6f
-- 
2.56.0

[PATCH 1/6] gdb/testsuite: make with_rocm_gpu_lock transparent to return

[Simon Marchi]

A "return" in the body of with_rocm_gpu_lock does not return from the
calling proc.  Instead, it only returns from with_rocm_gpu_lock, and
execution continues after the with_rocm_gpu_lock block.  For instance,
with:

  proc do_test {} {
      with_rocm_gpu_lock {
	  if {![runto_main]} {
	      return
	  }
	  ...
      }

      # Execution continues here if runto_main fails.
  }

Also, the "return" skips the gdb_exit call at the end of
with_rocm_gpu_lock, despite the comment saying that it is there
precisely for the case where the body returns early.  Same thing if the
body throws an error, or uses "break" or "continue".  Finally, gdb_exit
is called after releasing the GPU lock, while the comment above
with_rocm_gpu_lock says it is called before.

Before commit fbb0edfe60ed ("[gdb/testsuite] Factor out proc
with_lock", 2024-05-04), with_rocm_gpu_lock caught the result of the
body, called gdb_exit, released the lock and then propagated the result
to its caller.  That commit moved the catching and propagating to the
new with_lock proc.  Since then, a "return" in the body makes with_lock
return from with_rocm_gpu_lock, skipping gdb_exit, and the caller of
with_rocm_gpu_lock continues as if nothing happened.  Commit
7730d75c2a89 ("[gdb/testsuite] Refactor exception handling",
2026-09-23) later changed with_lock to use transparent_uplevel, but
kept the same behavior.

Fix it by using try/finally and transparent_uplevel inside the with_lock
block, like other wrapper procs in the testsuite.  The block passed to
with_lock is evaluated in the frame of with_rocm_gpu_lock, so
transparent_uplevel evaluates the body in the frame of the caller of
with_rocm_gpu_lock, and propagates control flow commands to it.  The
finally clause calls gdb_exit while still holding the lock.

Add gdb.testsuite/with-rocm-gpu-lock.exp, based on
gdb.testsuite/transparent-uplevel.exp, which checks that control flow
commands in the body of with_rocm_gpu_lock behave the same as if they
were not wrapped, and that GDB is exited in all cases.  I don't know if
we really need that test, it's not like we have a test for all wrapper
procs (we have a lot of them).  But I included it to demonstrate the
problems with the current with_rocm_gpu_lock, if someone wants to toy
with it.  I wouldn't mind dropping it from the patch.

Change-Id: Ie42db74105f03b08e5722edc5b5d910b3bc250b6
---
 .../gdb.testsuite/with-rocm-gpu-lock.exp      | 90 +++++++++++++++++++
 gdb/testsuite/lib/rocm.exp                    | 20 +++--
 2 files changed, 104 insertions(+), 6 deletions(-)
 create mode 100644 gdb/testsuite/gdb.testsuite/with-rocm-gpu-lock.exp

diff --git a/gdb/testsuite/gdb.testsuite/with-rocm-gpu-lock.exp b/gdb/testsuite/gdb.testsuite/with-rocm-gpu-lock.exp
new file mode 100644
index 000000000000..6bbe52e203e7
--- /dev/null
+++ b/gdb/testsuite/gdb.testsuite/with-rocm-gpu-lock.exp
@@ -0,0 +1,90 @@
+# Copyright 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/>.
+
+# Test that with_rocm_gpu_lock is transparent with respect to control flow
+# commands in its body, and that it exits GDB in all cases.
+
+load_lib rocm.exp
+
+set stmts \
+    [list \
+	 {} \
+	 "error foo" \
+	 return \
+	 "return foo" \
+	 break \
+	 continue \
+	 "return -code ok" \
+	 "return -code error" \
+	 "return -code return" \
+	 "return -code break" \
+	 "return -code continue" \
+	 "return -code ok -level 0" \
+	 "return -code ok -level 1" \
+	 "return -code ok -level 2" \
+	 "return -code error -level 0" \
+	 "return -code error -level 1" \
+	 "return -code error -level 2" \
+	 "return -code return -level 0" \
+	 "return -code return -level 1" \
+	 "return -code return -level 2" \
+	 "return -code break -level 0" \
+	 "return -code break -level 1" \
+	 "return -code break -level 2" \
+	 "return -code continue -level 0" \
+	 "return -code continue -level 1" \
+	 "return -code continue -level 2"]
+
+foreach stmt $stmts {
+    verbose -log "STMT: $stmt"
+
+    set res1 [catch {
+	{*}$stmt
+    } script_res1 opts1]
+
+    clean_restart
+
+    set res2 [catch {
+	with_rocm_gpu_lock {
+	    {*}$stmt
+	}
+    } script_res2 opts2]
+
+    gdb_assert {![info exists ::gdb_spawn_id]} "$stmt: gdb exited"
+
+    verbose -log "RES1: $res1"
+    verbose -log "RES2: $res2"
+    gdb_assert {$res1 == $res2} "$stmt: res"
+
+    verbose -log "SCRIPT_RES1: '$script_res1'"
+    verbose -log "SCRIPT_RES2: '$script_res2'"
+    gdb_assert {$script_res1 == $script_res2} "$stmt: script_res"
+
+    # Note: we skip -errorinfo and -errorstack here.  These differ for
+    # something like "error foo", and they should.  When it comes to errors,
+    # we're not trying to hide the indirection.
+    foreach key {-code -level -errorcode} {
+	set v1 {}
+	set v2 {}
+	if {[dict exists $opts1 $key]} {
+	    set v1 [dict get $opts1 $key]
+	}
+	if {[dict exists $opts2 $key]} {
+	    set v2 [dict get $opts2 $key]
+	}
+	verbose -log "KEY: $key, V1: $v1"
+	verbose -log "KEY: $key, V2: $v2"
+	gdb_assert {$v1 == $v2} "$stmt: opts: $key"
+    }
+}
diff --git a/gdb/testsuite/lib/rocm.exp b/gdb/testsuite/lib/rocm.exp
index f8a3b233faec..3a050769ab7c 100644
--- a/gdb/testsuite/lib/rocm.exp
+++ b/gdb/testsuite/lib/rocm.exp
@@ -162,15 +162,23 @@ gdb_caching_proc allow_hip_tests {} {
 # at a time.
 set gpu_lock_filename gpu-parallel.lock
 
-# Run body under the GPU lock.  Also calls gdb_exit before releasing
+# Run BODY under the GPU lock.  Also calls gdb_exit before releasing
 # the GPU lock.
+#
+# BODY is evaluated in the caller's frame, and control flow commands
+# in BODY behave as if BODY was not wrapped.  For instance, "return"
+# returns from the caller.
 
 proc with_rocm_gpu_lock { body } {
-    with_lock $::gpu_lock_filename {uplevel 1 $body}
-
-    # In case BODY returned early due to some testcase failing, and
-    # left GDB running, debugging the GPU.
-    gdb_exit
+    with_lock $::gpu_lock_filename {
+	try {
+	    transparent_uplevel $body
+	} finally {
+	    # In case BODY returned early due to some testcase failing, and
+	    # left GDB running, debugging the GPU.
+	    gdb_exit
+	}
+    }
 }
 
 # Return true if all the devices support debugging multiple processes
-- 
2.56.0

[PATCH 2/6] gdb/corefile: handle no signalled thread when writing NT_GDB_TDESC

[Simon Marchi]

Commit c14993e9dc5 ("gdb/corefile: write NT_GDB_TDESC based on
signalled thread") made linux_make_corefile_notes and
elf_none_make_corefile_notes use the signalled thread's architecture to
write the NT_GDB_TDESC note:

  gdbarch = target_thread_architecture (signalled_thr->ptid);

However, gcore_find_signalled_thread returns nullptr if the selected
thread has exited and no other thread was stopped by a signal.  The code
just above already checks for that, but this line does not, leading to
a null pointer dereference.

This can happen in non-stop mode, if the selected thread exits and the
other threads are then interrupted.  For instance, with an UBSan build:

  (gdb) thread
  [Current thread is 2 (LWP 301816) (exited)]
  (gdb) gcore
  .../gdb/linux-tdep.c:2586:40: runtime error: member access within null pointer of type 'struct thread_info'

(without UBSan it's just a segfault)

Fix this by falling back to the gdbarch passed to the function, which is
what was used before commit c14993e9dc5.

Add a test that reproduces the scenario above.

Change-Id: I058ec7a19dc7425e2728379317ba1dd879904eb9
---
 gdb/elf-none-tdep.c                           |  7 +-
 gdb/linux-tdep.c                              |  7 +-
 .../gcore-exited-selected-thread.c            | 51 ++++++++++++
 .../gcore-exited-selected-thread.exp          | 83 +++++++++++++++++++
 4 files changed, 144 insertions(+), 4 deletions(-)
 create mode 100644 gdb/testsuite/gdb.threads/gcore-exited-selected-thread.c
 create mode 100644 gdb/testsuite/gdb.threads/gcore-exited-selected-thread.exp

diff --git a/gdb/elf-none-tdep.c b/gdb/elf-none-tdep.c
index 0d063038f5ec..e046d99bf51b 100644
--- a/gdb/elf-none-tdep.c
+++ b/gdb/elf-none-tdep.c
@@ -113,8 +113,11 @@ elf_none_make_corefile_notes (struct gdbarch *gdbarch, bfd *obfd,
   /* Include the target description when possible.  Some architectures
      allow for per-thread gdbarch so we should really be emitting a tdesc
      per-thread, however, we don't currently support reading in a
-     per-thread tdesc, so just emit the tdesc for the signalled thread.  */
-  gdbarch = target_thread_architecture (signalled_thr->ptid);
+     per-thread tdesc, so just emit the tdesc for the signalled thread.  If
+     there is no signalled thread, fall back to GDBARCH.  */
+  if (signalled_thr != nullptr)
+    gdbarch = target_thread_architecture (signalled_thr->ptid);
+
   gcore_elf_make_tdesc_note (gdbarch, obfd, &note_data, note_size);
 
   return note_data;
diff --git a/gdb/linux-tdep.c b/gdb/linux-tdep.c
index 4f6910694d50..fbc96c0975a7 100644
--- a/gdb/linux-tdep.c
+++ b/gdb/linux-tdep.c
@@ -2582,8 +2582,11 @@ linux_make_corefile_notes (struct gdbarch *gdbarch, bfd *obfd, int *note_size)
   /* Include the target description when possible.  Some architectures
      allow for per-thread gdbarch so we should really be emitting a tdesc
      per-thread, however, we don't currently support reading in a
-     per-thread tdesc, so just emit the tdesc for the signalled thread.  */
-  gdbarch = target_thread_architecture (signalled_thr->ptid);
+     per-thread tdesc, so just emit the tdesc for the signalled thread.  If
+     there is no signalled thread, fall back to GDBARCH.  */
+  if (signalled_thr != nullptr)
+    gdbarch = target_thread_architecture (signalled_thr->ptid);
+
   gcore_elf_make_tdesc_note (gdbarch, obfd, &note_data, note_size);
 
   return note_data;
diff --git a/gdb/testsuite/gdb.threads/gcore-exited-selected-thread.c b/gdb/testsuite/gdb.threads/gcore-exited-selected-thread.c
new file mode 100644
index 000000000000..5a2fbf710831
--- /dev/null
+++ b/gdb/testsuite/gdb.threads/gcore-exited-selected-thread.c
@@ -0,0 +1,51 @@
+/* This testcase is part of GDB, the GNU debugger.
+
+   Copyright 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 <assert.h>
+#include <pthread.h>
+#include <unistd.h>
+
+static volatile int joined = 0;
+
+static void *
+thread_func (void *arg)
+{
+  return NULL; /* Break here.  */
+}
+
+int
+main (void)
+{
+  pthread_t thread;
+  int res;
+
+  alarm (30);
+
+  res = pthread_create (&thread, NULL, thread_func, NULL);
+  assert (res == 0);
+
+  res = pthread_join (thread, NULL);
+  assert (res == 0);
+
+  joined = 1; /* After join.  */
+
+  /* Keep running until GDB interrupts us.  */
+  while (1)
+    sleep (1);
+
+  return 0;
+}
diff --git a/gdb/testsuite/gdb.threads/gcore-exited-selected-thread.exp b/gdb/testsuite/gdb.threads/gcore-exited-selected-thread.exp
new file mode 100644
index 000000000000..10a89ffa9060
--- /dev/null
+++ b/gdb/testsuite/gdb.threads/gcore-exited-selected-thread.exp
@@ -0,0 +1,83 @@
+# Copyright 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/>.
+
+# Test generating a core file when the selected thread has exited and
+# no other thread was stopped by a signal.  In that case, there is no
+# "signalled thread" to dump first.
+
+standard_testfile
+
+if { [build_executable "failed to prepare" $testfile $srcfile \
+	  {debug pthreads}] == -1 } {
+    return
+}
+
+save_vars { ::GDBFLAGS } {
+    append ::GDBFLAGS " -ex \"set non-stop on\""
+    clean_restart $::testfile
+}
+
+gdb_breakpoint [gdb_get_line_number "Break here."]
+gdb_run_cmd
+gdb_test "" "Thread 2 \[^\r\n\]* hit Breakpoint $decimal, thread_func .*" \
+    "run to breakpoint in thread"
+gdb_test "thread 2" "Switching to thread 2 .*" "select thread 2"
+
+# Print a message when the main thread has joined thread 2.  This tells
+# us that thread 2 has exited, without relying on GDB reporting the
+# thread exit right away, which it doesn't do with remote targets.  Use
+# a dprintf rather than a breakpoint, so that the main thread does not
+# stop with a SIGTRAP and become the signalled thread.
+gdb_test "dprintf [gdb_get_line_number "After join."],\"thread 2 joined\\n\"" \
+    "Dprintf $decimal at .*" "set dprintf after join"
+
+# Resume thread 2 and let it exit.  It stays the selected thread.
+gdb_test_multiple "continue &" "continue thread 2 until exit" {
+    -re "Continuing\\.\r\n$gdb_prompt " {
+	exp_continue
+    }
+    -re "thread 2 joined\r\n" {
+	pass $gdb_test_name
+    }
+}
+
+# Stop the main thread.  Being interrupted, it doesn't have a stop
+# signal.
+gdb_test_multiple "interrupt -a" "" {
+    -re "^interrupt -a\r\n$gdb_prompt " {
+	exp_continue
+    }
+    -re "Thread 1 \[^\r\n\]* stopped\\.\r\n" {
+	pass $gdb_test_name
+    }
+}
+
+gdb_test "thread" "Current thread is 2 \\(.*\\) \\(exited\\)\\\]" \
+    "selected thread has exited"
+
+set corefile [host_standard_output_file $testfile.core]
+if { [gdb_gcore_cmd $corefile "save a corefile"] != 1 } {
+    return
+}
+
+clean_restart $::testfile
+if { [gdb_core_cmd $corefile "load corefile"] != 1 } {
+    return
+}
+
+# Check that the core file contains only the main thread.
+gdb_test "info threads" \
+    "Id +Target Id +Frame *\r\n\\* 1 +Thread \[^\r\n\]*" \
+    "core file contains main thread"
-- 
2.56.0

[PATCH 3/6] gdb/amd-dbgapi: handle regno == -1 in fetch_registers and store_registers

[Simon Marchi]

Running "gcore" while a GPU wave is stopped at a breakpoint makes GDB
crash:

  $ gdb -q -nx testsuite/outputs/gdb.rocm/simple/simple \
      -ex "with breakpoint pending on -- break do_an_addition" \
      -ex r -ex gcore -batch
  ...
  regcache.c:162: internal-error: register_type: Assertion `regnum >= 0 && regnum < descr->nr_cooked_registers' failed.

The backtrace is:

    #0  internal_error_loc (file=0x55556355d500 "/home/smarchi/src/binutils-gdb/gdb/regcache.c", line=162, fmt=0x55556355cf60 "%s: Assertion `%s' failed.") at /home/smarchi/src/binutils-gdb/gdbsupport/errors.cc:53
    #1  register_type (gdbarch=0x51c000073880, regnum=-1) at /home/smarchi/src/binutils-gdb/gdb/regcache.c:162
    #2  amd_dbgapi_target::fetch_registers (this=0x555568df2760 <the_amd_dbgapi_target>, regcache=0x5070000c9420, regno=-1) at /home/smarchi/src/binutils-gdb/gdb/amd-dbgapi-target.c:2039
    #3  target_fetch_registers (regcache=0x5070000c9420, regno=-1) at /home/smarchi/src/binutils-gdb/gdb/target.c:3925
    #4  gcore_elf_build_thread_register_notes (gdbarch=0x51c000073880, info=0x51800008b480, stop_signal=GDB_SIGNAL_TRAP, obfd=0x5120001c0b40, note_data=0x7fffe698d420, note_size=0x7fffe68c0e30) at /home/smarchi/src/binutils-gdb/gdb/gcore-elf.c:132
    #5  linux_corefile_thread (info=0x51800008b480, gdbarch=0x51c000073880, obfd=0x5120001c0b40, note_data=std::unique_ptr<char> = {...}, note_size=0x7fffe68c0e30, stop_signal=GDB_SIGNAL_TRAP) at /home/smarchi/src/binutils-gdb/gdb/linux-tdep.c:1957
    #6  linux_make_corefile_notes (gdbarch=0x51c00002d080, obfd=0x5120001c0b40, note_size=0x7fffe68c0e30) at /home/smarchi/src/binutils-gdb/gdb/linux-tdep.c:2543
    #7  gdbarch_make_corefile_notes (gdbarch=0x51c00002d080, obfd=0x5120001c0b40, note_size=0x7fffe68c0e30) at /home/smarchi/src/binutils-gdb/gdb/gdbarch-gen.c:3664
    #8  write_gcore_file_1 (obfd=0x5120001c0b40) at /home/smarchi/src/binutils-gdb/gdb/gcore.c:95
    #9  write_gcore_file (obfd=0x5120001c0b40) at /home/smarchi/src/binutils-gdb/gdb/gcore.c:137
    #10 gcore_command (args=0x0, from_tty=0) at /home/smarchi/src/binutils-gdb/gdb/gcore.c:175

gcore_elf_build_thread_register_notes calls target_fetch_registers with
regno == -1, meaning "fetch all registers".  The amd-dbgapi target
assumes regno is always a valid register number, and passes -1 directly
to register_type.  store_registers has the same problem.

Make both functions handle regno == -1 by iterating over all raw
registers.  This code is mostly taken from ROCgdb, except that it also
emits the warning when some registers fail to read or write in the
"regno == -1" cases.  ROCgdb only warns in the single register case.

Like ROCgdb, invalidate volatile registers once, after all registers are
written, so that writing a register with the "invalidate volatile"
property (e.g. VCC) does not invalidate a volatile register (e.g.
STATUS) that has yet to be written.

With this fix, the gcore command above gets a bit further, but now hits:

  gdbarch-gen.c:3637: internal-error: gdbarch_iterate_over_regset_sections: Assertion `gdbarch->iterate_over_regset_sections != nullptr' failed.

This is because linux_make_corefile_notes tries to dump GPU waves as if
they were host threads, using the amdgpu gdbarch, which does not
implement iterate_over_regset_sections.  This is addressed by a
following patch.

Change-Id: Ie669eb8a538ddfd3f35745e3af53eaf5541aa01b
---
 gdb/amd-dbgapi-target.c | 100 +++++++++++++++++++++++++---------------
 1 file changed, 64 insertions(+), 36 deletions(-)

diff --git a/gdb/amd-dbgapi-target.c b/gdb/amd-dbgapi-target.c
index 2fd175de6f7d..099deceaabad 100644
--- a/gdb/amd-dbgapi-target.c
+++ b/gdb/amd-dbgapi-target.c
@@ -2035,20 +2035,32 @@ amd_dbgapi_target::fetch_registers (struct regcache *regcache, int regno)
 
   amdgpu_gdbarch_tdep *tdep = get_amdgpu_gdbarch_tdep (gdbarch);
   amd_dbgapi_wave_id_t wave_id = get_amd_dbgapi_wave_id (regcache->ptid ());
-  gdb_byte raw[AMDGPU_MAX_REGISTER_SIZE];
-  ULONGEST reg_size = register_type (gdbarch, regno)->length ();
-  gdb_assert (reg_size <= AMDGPU_MAX_REGISTER_SIZE);
 
-  amd_dbgapi_status_t status
-    = amd_dbgapi_read_register (wave_id, tdep->register_ids[regno], 0,
-				reg_size, raw);
+  auto fetch_one_register = [&] (int r)
+    {
+      gdb_byte raw[AMDGPU_MAX_REGISTER_SIZE];
+      ULONGEST reg_size = register_type (gdbarch, r)->length ();
+      gdb_assert (reg_size <= AMDGPU_MAX_REGISTER_SIZE);
 
-  if (status == AMD_DBGAPI_STATUS_SUCCESS)
-    regcache->raw_supply (regno, raw);
-  else if (status != AMD_DBGAPI_STATUS_ERROR_REGISTER_NOT_AVAILABLE)
-    warning (_("Couldn't read register %s (#%d) (%s)."),
-	     gdbarch_register_name (gdbarch, regno), regno,
-	     get_status_string (status));
+      amd_dbgapi_status_t status
+	= amd_dbgapi_read_register (wave_id, tdep->register_ids[r], 0,
+				    reg_size, raw);
+
+      if (status == AMD_DBGAPI_STATUS_SUCCESS)
+	regcache->raw_supply (r, raw);
+      else if (status != AMD_DBGAPI_STATUS_ERROR_REGISTER_NOT_AVAILABLE)
+	warning (_("Couldn't read register %s (#%d) (%s)."),
+		 gdbarch_register_name (gdbarch, r), r,
+		 get_status_string (status));
+    };
+
+  if (regno == -1)
+    {
+      for (int r = 0; r < gdbarch_num_regs (gdbarch); ++r)
+	fetch_one_register (r);
+    }
+  else
+    fetch_one_register (regno);
 }
 
 void
@@ -2063,38 +2075,54 @@ amd_dbgapi_target::store_registers (struct regcache *regcache, int regno)
   struct gdbarch *gdbarch = regcache->arch ();
   gdb_assert (is_amdgpu_arch (gdbarch));
 
-  gdb_byte raw[AMDGPU_MAX_REGISTER_SIZE];
-  ULONGEST reg_size = register_type (gdbarch, regno)->length ();
-  gdb_assert (reg_size <= AMDGPU_MAX_REGISTER_SIZE);
-  regcache->raw_collect (regno, &raw);
-
   amdgpu_gdbarch_tdep *tdep = get_amdgpu_gdbarch_tdep (gdbarch);
+  amd_dbgapi_wave_id_t wave_id = get_amd_dbgapi_wave_id (regcache->ptid ());
+  bool invalidate_volatile = false;
 
-  /* If the register has read-only bits, invalidate the value in the regcache
-     as the value actually written may differ.  */
-  if (tdep->register_properties[regno]
-      & AMD_DBGAPI_REGISTER_PROPERTY_READONLY_BITS)
-    regcache->invalidate (regno);
+  auto store_one_register = [&] (int r)
+    {
+      gdb_byte raw[AMDGPU_MAX_REGISTER_SIZE];
+      ULONGEST reg_size = register_type (gdbarch, r)->length ();
+      gdb_assert (reg_size <= AMDGPU_MAX_REGISTER_SIZE);
+      regcache->raw_collect (r, &raw);
 
-  /* Invalidate all volatile registers if this register has the invalidate
-     volatile property.  For example, writing to VCC may change the content
-     of STATUS.VCCZ.  */
-  if (tdep->register_properties[regno]
-      & AMD_DBGAPI_REGISTER_PROPERTY_INVALIDATE_VOLATILE)
+      /* If the register has read-only bits, invalidate the value in the
+	 regcache as the value actually written may differ.  */
+      if (tdep->register_properties[r]
+	  & AMD_DBGAPI_REGISTER_PROPERTY_READONLY_BITS)
+	regcache->invalidate (r);
+
+      /* Writing this register may change the value of volatile registers.
+	 For example, writing to VCC may change the content of STATUS.VCCZ.
+	 Invalidate them once all registers are written, so that we don't
+	 invalidate a register we have yet to write.  */
+      if (tdep->register_properties[r]
+	  & AMD_DBGAPI_REGISTER_PROPERTY_INVALIDATE_VOLATILE)
+	invalidate_volatile = true;
+
+      amd_dbgapi_status_t status
+	= amd_dbgapi_write_register (wave_id, tdep->register_ids[r], 0,
+				     reg_size, raw);
+
+      if (status != AMD_DBGAPI_STATUS_SUCCESS)
+	warning (_("Couldn't write register %s (#%d)."),
+		 gdbarch_register_name (gdbarch, r), r);
+    };
+
+  if (regno == -1)
+    {
+      for (int r = 0; r < gdbarch_num_regs (gdbarch); ++r)
+	store_one_register (r);
+    }
+  else
+    store_one_register (regno);
+
+  if (invalidate_volatile)
     {
       for (size_t r = 0; r < tdep->register_properties.size (); ++r)
 	if (tdep->register_properties[r] & AMD_DBGAPI_REGISTER_PROPERTY_VOLATILE)
 	  regcache->invalidate (r);
     }
-
-  amd_dbgapi_wave_id_t wave_id = get_amd_dbgapi_wave_id (regcache->ptid ());
-  amd_dbgapi_status_t status
-    = amd_dbgapi_write_register (wave_id, tdep->register_ids[regno], 0,
-				 reg_size, raw);
-
-  if (status != AMD_DBGAPI_STATUS_SUCCESS)
-    warning (_("Couldn't write register %s (#%d)."),
-	     gdbarch_register_name (gdbarch, regno), regno);
 }
 
 struct gdbarch *
-- 
2.56.0

[PATCH 4/6] gdb/linux-tdep: don't write core file notes for non-Linux threads

[Simon Marchi]

Running "gcore" while debugging a ROCm program, with some GPU waves
stopped, makes GDB crash:

  (gdb) gcore
  .../gdb/gdbarch-gen.c:3637: internal-error: gdbarch_iterate_over_regset_sections: Assertion `gdbarch->iterate_over_regset_sections != nullptr' failed.

linux_make_corefile_notes writes register notes for every thread of the
inferior, using each thread's own gdbarch.  When debugging a ROCm
program, the inferior's thread list contains both the host (Linux)
threads and the GPU waves, provided by the amd-dbgapi target.  The waves
use the amdgcn gdbarch, which doesn't implement
iterate_over_regset_sections, hence the assertion failure.  This happens
regardless of which thread is selected, since all threads are visited.

There is also a related problem: if a wave is the signalled thread, for
instance because it is the selected thread and hit a breakpoint, it is
dumped first, its stop signal is used for all the other threads, and its
gdbarch is used to write the NT_GDB_TDESC note.  This would produce a
core file with an amdgcn target description.

Make linux_make_corefile_notes skip non-Linux threads, and only consider
Linux threads when looking for the signalled thread.

To identify Linux threads, add the is_linux_thread function.  It relies
on the fact that the amd-dbgapi target uses ptids of the form (pid, 1,
wave_id), with a non-zero tid field, while Linux threads have a zero tid
field.  This is not super pretty, but we need to teach the Linux support
that there might be non-Linux threads in the thread list one way or
another, and this is cheap, simple and reliable.  If we ever find a
better way to discriminate threads, then we just have to update this
function.

Add an optional filter parameter to gcore_find_signalled_thread, to
let linux_make_corefile_notes exclude non-Linux threads.

Add the gdb.rocm/gcore-host.exp test, which generates a core file while
a GPU wave is stopped at a breakpoint, and a host thread is selected.
It then loads the core file and checks that it contains all the host
threads, and only them.  Generating a core file while a GPU wave is
selected still fails, because of a separate bug that is addressed in a
following patch.

The test is named like that to avoid clashing with the
gdb.rocm/gcore.exp test downstream, which tests core generation
including the GPU threads.  This feature is going to be upstreamed
eventually, but in the meantime the gdb.rocm/gcore-host.exp test is
meant to test that the presence of GPU waves does not prevent a
host-only core dump from being created.

Change-Id: I33fd7fcf21ca8c645195a821e8dc65f042cda364
---
 gdb/gcore.c                           |  27 ++++---
 gdb/gcore.h                           |   9 ++-
 gdb/linux-tdep.c                      |  29 +++++++-
 gdb/testsuite/gdb.rocm/gcore-host.cpp |  42 +++++++++++
 gdb/testsuite/gdb.rocm/gcore-host.exp | 102 ++++++++++++++++++++++++++
 5 files changed, 195 insertions(+), 14 deletions(-)
 create mode 100644 gdb/testsuite/gdb.rocm/gcore-host.cpp
 create mode 100644 gdb/testsuite/gdb.rocm/gcore-host.exp

diff --git a/gdb/gcore.c b/gdb/gcore.c
index 77d915505d1c..1336e56cda82 100644
--- a/gdb/gcore.c
+++ b/gdb/gcore.c
@@ -903,22 +903,31 @@ gcore_memory_sections (bfd *obfd)
 /* See gcore.h.  */
 
 thread_info *
-gcore_find_signalled_thread ()
+gcore_find_signalled_thread
+  (gdb::function_view<bool (const thread_info &)> filter)
 {
+  auto accept = [filter] (const thread_info &thr)
+    {
+      return filter == nullptr || filter (thr);
+    };
+
   thread_info *curr_thr = inferior_thread ();
-  if (curr_thr->state () != THREAD_EXITED
-      && curr_thr->stop_signal () != GDB_SIGNAL_0)
+
+  /* Give priority to the current thread, unless it has exited or is not
+     accepted by FILTER.  */
+  if (curr_thr->state () == THREAD_EXITED || !accept (*curr_thr))
+    curr_thr = nullptr;
+
+  if (curr_thr != nullptr && curr_thr->stop_signal () != GDB_SIGNAL_0)
     return curr_thr;
 
   for (thread_info &thr : current_inferior ()->non_exited_threads ())
-    if (thr.stop_signal () != GDB_SIGNAL_0)
+    if (accept (thr) && thr.stop_signal () != GDB_SIGNAL_0)
       return &thr;
 
-  /* Default to the current thread, unless it has exited.  */
-  if (curr_thr->state () != THREAD_EXITED)
-    return curr_thr;
-
-  return nullptr;
+  /* Default to the current thread, unless it has exited or is not
+     accepted by FILTER, in which case return nullptr.  */
+  return curr_thr;
 }
 
 INIT_GDB_FILE (gcore)
diff --git a/gdb/gcore.h b/gdb/gcore.h
index 2431e123b60e..277c28d9c9ac 100644
--- a/gdb/gcore.h
+++ b/gdb/gcore.h
@@ -22,6 +22,7 @@
 
 #include "find-memory-region.h"
 #include "gdb_bfd.h"
+#include "gdbsupport/function-view.h"
 
 struct thread_info;
 
@@ -33,8 +34,12 @@ extern bool objfile_find_memory_regions (struct target_ops *self,
 /* Find the signalled thread.  In case there's more than one signalled
    thread, prefer the current thread, if it is signalled.  If no thread was
    signalled, default to the current thread, unless it has exited, in which
-   case return NULL.  */
+   case return nullptr.
 
-extern thread_info *gcore_find_signalled_thread ();
+   If FILTER is non-null, only consider threads for which FILTER returns
+   true.  */
+
+extern thread_info *gcore_find_signalled_thread
+  (gdb::function_view<bool (const thread_info &)> filter = nullptr);
 
 #endif /* GDB_GCORE_H */
diff --git a/gdb/linux-tdep.c b/gdb/linux-tdep.c
index fbc96c0975a7..3491c63c017a 100644
--- a/gdb/linux-tdep.c
+++ b/gdb/linux-tdep.c
@@ -456,6 +456,26 @@ linux_has_shared_address_space (struct gdbarch *gdbarch)
   return linux_is_uclinux ();
 }
 
+/* Return true if THREAD appears to be a Linux thread.
+
+   When debugging a ROCm program, an inferior contains both host threads
+   (provided by the linux-nat target) and GPU threads (provided by the
+   amd-dbgapi target).
+
+   To differentiate them from the host threads, the amd-dbgapi target uses
+   ptids of the form
+
+     (pid, 1, wave_id)
+
+   where wave_id is a non-zero unique number.  Return true (is a Linux
+   thread) if the tid field is 0.  */
+
+static bool
+is_linux_thread (const thread_info &thread)
+{
+  return thread.ptid.tid () == 0;
+}
+
 /* Return the ptid of a thread of the current inferior whose lwp id can be
    used to build "/proc/<lwp>" paths.
 
@@ -2524,8 +2544,11 @@ linux_make_corefile_notes (struct gdbarch *gdbarch, bfd *obfd, int *note_size)
 
   /* Like the kernel, prefer dumping the signalled thread first.
      "First thread" is what tools use to infer the signalled
-     thread.  */
-  thread_info *signalled_thr = gcore_find_signalled_thread ();
+     thread.
+
+     Only consider Linux threads.  Other threads, like AMD GPU waves, can't be
+     described using Linux thread notes.  */
+  thread_info *signalled_thr = gcore_find_signalled_thread (is_linux_thread);
   gdb_signal stop_signal;
   if (signalled_thr != nullptr)
     stop_signal = signalled_thr->stop_signal ();
@@ -2546,7 +2569,7 @@ linux_make_corefile_notes (struct gdbarch *gdbarch, bfd *obfd, int *note_size)
     }
   for (thread_info &thr : current_inferior ()->non_exited_threads ())
     {
-      if (&thr == signalled_thr)
+      if (&thr == signalled_thr || !is_linux_thread (thr))
 	continue;
 
       /* On some architectures, like AArch64, each thread can have a distinct
diff --git a/gdb/testsuite/gdb.rocm/gcore-host.cpp b/gdb/testsuite/gdb.rocm/gcore-host.cpp
new file mode 100644
index 000000000000..0f8ca14fe90f
--- /dev/null
+++ b/gdb/testsuite/gdb.rocm/gcore-host.cpp
@@ -0,0 +1,42 @@
+/* This testcase is part of GDB, the GNU debugger.
+
+   Copyright 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 "hip/hip_runtime.h"
+#include <cassert>
+
+__global__ void
+kernel (int *out)
+{
+  *out = 1;
+}
+
+int
+main ()
+{
+  int *result_ptr, result;
+
+  hipError_t error = hipMalloc (&result_ptr, sizeof (int));
+  assert (error == hipSuccess);
+
+  kernel<<<dim3 (1), dim3 (1), 0, 0>>> (result_ptr);
+
+  error = hipMemcpyDtoH (&result, result_ptr, sizeof (int));
+  assert (error == hipSuccess);
+  assert (result == 1);
+
+  return 0;
+}
diff --git a/gdb/testsuite/gdb.rocm/gcore-host.exp b/gdb/testsuite/gdb.rocm/gcore-host.exp
new file mode 100644
index 000000000000..80bf60e4e9f5
--- /dev/null
+++ b/gdb/testsuite/gdb.rocm/gcore-host.exp
@@ -0,0 +1,102 @@
+# Copyright 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/>.
+
+# Test generating a core file for an inferior that has GPU waves, then
+# loading it.  GDB does not support writing GPU waves to core files, so the
+# core file is expected to contain only the host threads.
+
+load_lib rocm.exp
+
+standard_testfile .cpp
+
+require allow_hip_tests
+
+if {[build_executable "failed to prepare" $testfile $srcfile {debug hip}]} {
+    return
+}
+
+# Generate a core file while a GPU wave is stopped at a breakpoint, with
+# SELECTED_THREAD ("host" or "wave") being the selected thread.  Then load
+# the core file and verify its contents.
+
+proc do_test { selected_thread } {
+    clean_restart $::testfile
+
+    set corefile [standard_output_file $::testfile.$selected_thread.core]
+
+    with_rocm_gpu_lock {
+	if {![runto_main]} {
+	    return
+	}
+
+	gdb_breakpoint "kernel" -allow-pending
+
+	gdb_test "continue" \
+	    "Thread $::decimal hit Breakpoint ${::decimal}(?:\\.${::decimal})?, kernel .*"
+
+	# Count the host threads, for comparison with the core file.
+	set num_host_threads 0
+	set num_waves 0
+	gdb_test_multiple "info threads" "" -lbl {
+	    -re "\r\n\\*? +$::decimal +Thread $::hex \\(LWP $::decimal\\)\[^\r\n\]*" {
+		incr num_host_threads
+		exp_continue
+	    }
+	    -re "\r\n\\*? +$::decimal +AMDGPU Wave\[^\r\n\]*" {
+		incr num_waves
+		exp_continue
+	    }
+	    -re -wrap "" {
+		gdb_assert { $num_host_threads > 0 && $num_waves > 0 } \
+		    $gdb_test_name
+	    }
+	}
+
+	if { $selected_thread == "host" } {
+	    gdb_test "thread 1" "Switching to thread 1 .*"
+	}
+
+	if { [gdb_gcore_cmd $corefile "save a corefile"] != 1 } {
+	    return
+	}
+    }
+
+    clean_restart $::testfile
+
+    gdb_test "core-file $corefile" \
+	"Core was generated by `[string_to_regexp $::binfile]'\\..*" \
+	"load corefile"
+
+    # The core file is expected to contain all the host threads, and nothing
+    # else.
+    set num_core_threads 0
+    gdb_test_multiple "info threads" "core file contains host threads" -lbl {
+	-re "\r\n\\*? +$::decimal +Thread $::hex \\(LWP $::decimal\\)\[^\r\n\]*" {
+	    incr num_core_threads
+	    exp_continue
+	}
+	-re "\r\n\\*? +$::decimal +\[^\r\n\]*" {
+	    fail $gdb_test_name
+	}
+	-re -wrap "" {
+	    gdb_assert { $num_core_threads == $num_host_threads } \
+		$gdb_test_name
+	}
+    }
+}
+
+foreach_with_prefix selected_thread { host } {
+    do_test $selected_thread
+}
-- 
2.56.0

[PATCH 5/6] gdb/amd-dbgapi: defer TARGET_OBJECT_AUXV xfers to the beneath target

[Simon Marchi]

From: Shahab Vahedi <shahab.vahedi@amd.com>

Reading the auxiliary vector of the inferior fails when a GPU wave is
the selected thread.  The read goes through
amd_dbgapi_target::xfer_partial, which only handles TARGET_OBJECT_MEMORY
for GPU threads, and returns TARGET_XFER_E_IO for everything else.

Most users of the auxiliary vector, including "info auxv", use the copy
cached in the inferior, which is normally populated early, while a host
thread is selected.  But if the cache is cleared while a wave is
selected, for instance by the "file" command, the next read fails, and
the failure is cached too.

A victim of that bug is the "gcore" command.  linux_make_corefile_notes
reads the auxiliary vector directly from the target (bypassing the
cache) to write the NT_AUXV note, so the note is always missing when a
GPU wave is selected.  As a consequence, when loading such a core file,
GDB can't find the shared libraries of a position independent
executable, and can't find the full program name (AT_EXECFN), so it
falls back to the possibly truncated name from the NT_PRPSINFO note.

Fix this by making amd_dbgapi_target::xfer_partial always let the target
beneath handle TARGET_OBJECT_AUXV.

Add the gdb.rocm/auxv-wave-selected.exp test, which reproduces the
"file" scenario above (perhaps we should have a maintenance command to
flush the auxv cache though).  The gcore scenario can't be tested yet,
because generating a core file while a GPU wave is selected fails for
another reason, fixed in the following patch.

The amd-dbgapi-target.c part of this change comes from the downstream
ROCgdb port, commit 47716f7de58 ("gdb/amd-dbgapi: defer
TARGET_OBJECT_AUXV xfers to the beneath target").

Co-Authored-By: Simon Marchi <simon.marchi@efficios.com>
Change-Id: I177eba5780706e4065e1b17f370a17dbea016d8d
---
 gdb/amd-dbgapi-target.c                       |  9 ++-
 gdb/testsuite/gdb.rocm/auxv-wave-selected.cpp | 42 ++++++++++++
 gdb/testsuite/gdb.rocm/auxv-wave-selected.exp | 65 +++++++++++++++++++
 3 files changed, 114 insertions(+), 2 deletions(-)
 create mode 100644 gdb/testsuite/gdb.rocm/auxv-wave-selected.cpp
 create mode 100644 gdb/testsuite/gdb.rocm/auxv-wave-selected.exp

diff --git a/gdb/amd-dbgapi-target.c b/gdb/amd-dbgapi-target.c
index 099deceaabad..1c6edb5a7fa3 100644
--- a/gdb/amd-dbgapi-target.c
+++ b/gdb/amd-dbgapi-target.c
@@ -752,9 +752,14 @@ amd_dbgapi_target::xfer_partial (enum target_object object, const char *annex,
      the parent, and inferior_ptid has the child's PTID.  We can use
      the parent's process_stratum_target (the one below us) to do the
      memory operation, but dbgapi knows nothing about the child and
-     would try to update the parent's memory.  */
+     would try to update the parent's memory.
+
+     The auxiliary vector is a property of the host process, so let the
+     target beneath handle it, even if the current thread is a GPU
+     thread.  */
   if ((!ptid_is_gpu (inferior_ptid) && object != TARGET_OBJECT_MEMORY)
-      || inferior_ptid.pid () != current_inferior ()->pid)
+      || inferior_ptid.pid () != current_inferior ()->pid
+      || object == TARGET_OBJECT_AUXV)
     return beneath ()->xfer_partial (object, annex, readbuf, writebuf,
 				     offset, requested_len, xfered_len);
 
diff --git a/gdb/testsuite/gdb.rocm/auxv-wave-selected.cpp b/gdb/testsuite/gdb.rocm/auxv-wave-selected.cpp
new file mode 100644
index 000000000000..0f8ca14fe90f
--- /dev/null
+++ b/gdb/testsuite/gdb.rocm/auxv-wave-selected.cpp
@@ -0,0 +1,42 @@
+/* This testcase is part of GDB, the GNU debugger.
+
+   Copyright 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 "hip/hip_runtime.h"
+#include <cassert>
+
+__global__ void
+kernel (int *out)
+{
+  *out = 1;
+}
+
+int
+main ()
+{
+  int *result_ptr, result;
+
+  hipError_t error = hipMalloc (&result_ptr, sizeof (int));
+  assert (error == hipSuccess);
+
+  kernel<<<dim3 (1), dim3 (1), 0, 0>>> (result_ptr);
+
+  error = hipMemcpyDtoH (&result, result_ptr, sizeof (int));
+  assert (error == hipSuccess);
+  assert (result == 1);
+
+  return 0;
+}
diff --git a/gdb/testsuite/gdb.rocm/auxv-wave-selected.exp b/gdb/testsuite/gdb.rocm/auxv-wave-selected.exp
new file mode 100644
index 000000000000..637c5f2561b6
--- /dev/null
+++ b/gdb/testsuite/gdb.rocm/auxv-wave-selected.exp
@@ -0,0 +1,65 @@
+# Copyright 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/>.
+
+# Test that GDB can read the auxiliary vector of the inferior while a GPU
+# wave is selected.
+
+load_lib rocm.exp
+
+standard_testfile .cpp
+
+require allow_hip_tests
+
+# Build a position-independent executable.  To find the shared libraries
+# loaded by a PIE, GDB needs the auxiliary vector to find where the
+# executable is loaded.
+if {[build_executable "failed to prepare" $testfile $srcfile {debug hip pie}]} {
+    return
+}
+
+clean_restart $testfile
+
+with_rocm_gpu_lock {
+    if {![runto_main]} {
+	return
+    }
+
+    gdb_breakpoint "kernel" -allow-pending
+
+    gdb_test "continue" \
+	"Thread $::decimal hit Breakpoint ${::decimal}(?:\\.${::decimal})?, kernel .*"
+
+    gdb_test "thread" "Current thread is $::decimal \\(AMDGPU Wave .*" \
+	"wave is selected"
+
+    delete_breakpoints
+
+    # GDB caches the auxiliary vector, and the cache is normally populated
+    # early, while a host thread is selected.  Use the "file" command while a
+    # wave is selected, which clears the cache, so that the auxiliary vector is
+    # read again while the wave is selected.
+    gdb_test "with confirm off -- file $::binfile" \
+	"Reading symbols from [string_to_regexp $::binfile]\\.\\.\\..*" \
+	"reload executable"
+
+    gdb_test "info auxv" \
+	"AT_ENTRY +Entry point of program +$::hex\r\n.*AT_NULL +End of vector +0x0" \
+	"info auxv with wave selected"
+
+    # Check that GDB still finds the shared libraries loaded by the
+    # program.
+    gdb_test "info sharedlibrary" "/libamdhip64\\.so.*" \
+	"shared libraries after reload"
+}
-- 
2.56.0

[PATCH 6/6] gdb/linux-tdep: ignore non-Linux threads in get_ptid_for_slash_proc

[Simon Marchi]

Commit 2f4e6ea8c05 ("gdb: use the selected thread's LWP ID when reading
Linux procfs files") made get_ptid_for_slash_proc prefer the selected
thread.  It does not consider that the inferior's thread list may
contain threads that are not Linux threads, such as AMD GPU waves.  The
amd-dbgapi target numbers waves like this:

  (pid, 1, wave_id)

where pid is the host process' pid and wave_id a non-zero number handed
out by the amd-dbgapi library.  If a GPU wave is selected,
get_ptid_for_slash_proc returns its ptid, and GDB attempts to read files
under "/proc/1/", which is bad.

This happens when running "gcore" while a GPU wave is selected.  As a
result, a bunch of information coming for /proc is missing from the core
file (unless you are running GDB as root, in which case it successfully
reads the information for the init process, but then it's the wrong
info).

Fix this by making get_ptid_for_slash_proc aware that non-Linux threads
may exist, using the is_linux_thread function added by a previous patch.
The logic is now to use the selected thread if it has not exited and it
is a Linux thread.  Otherwise, look for any non-exited Linux thread.

Extend gdb.rocm/gcore-host.exp to also generate a core file while a wave is
selected.  The test also covers the fix from the previous patch, which makes the
NT_AUXV note present in that case.  To verify it:

 - Check that "info auxv" shows the auxiliary vector when the core file is
   loaded.

 - Run the program with a long argument, and check that GDB shows the full
   command line when loading the core file.  GDB reads the command line from
   memory, starting at the address given by AT_EXECFN in the auxiliary vector.
   Without the auxiliary vector, it falls back to the command line saved in the
   NT_PRPSINFO note, which is truncated to 80 characters.

 - Build the test program as a PIE, since loading the core file of a PIE
   requires the auxiliary vector to find where the executable and the shared
   libraries are loaded, and check that "info sharedlibrary" lists the HIP
   runtime library.

Change-Id: I3582344e25fa632bb21d77b38bd3b755c6cc644b
---
 gdb/linux-tdep.c                      | 28 ++++++++++++++++----------
 gdb/testsuite/gdb.rocm/gcore-host.exp | 29 ++++++++++++++++++++++++---
 2 files changed, 43 insertions(+), 14 deletions(-)

diff --git a/gdb/linux-tdep.c b/gdb/linux-tdep.c
index 3491c63c017a..d42195c4ebe7 100644
--- a/gdb/linux-tdep.c
+++ b/gdb/linux-tdep.c
@@ -479,23 +479,29 @@ is_linux_thread (const thread_info &thread)
 /* Return the ptid of a thread of the current inferior whose lwp id can be
    used to build "/proc/<lwp>" paths.
 
-   Through any_non_exited_thread_of_inferior, this function prefers the
-   currently selected thread, if it has not exited.  For most information in
-   /proc, it does not matter which thread of a process we pick.  But a few
-   entries (like stat counters) are thread-specific.  This allows the user
-   to focus a particular thread and get specific information about that
-   thread.
+   This function prefers the currently selected thread, if it has not exited
+   and is a Linux thread.  For most information in /proc, it does not matter
+   which thread of a process we pick.  But a few entries (like stat counters)
+   are thread-specific.  This allows the user to focus a particular thread
+   and get specific information about that thread.
 
-   Throw an error if the inferior has no thread that GDB believes is
+   Throw an error if the inferior has no Linux thread that GDB believes is
    alive.  */
 
 static ptid_t
 get_ptid_for_slash_proc ()
 {
-  if (thread_info *thread
-	= any_non_exited_thread_of_inferior (current_inferior ());
-      thread != nullptr)
-    return thread->ptid;
+  if (inferior_ptid != null_ptid)
+    {
+      thread_info *thread = inferior_thread ();
+
+      if (thread->state () != THREAD_EXITED && is_linux_thread (*thread))
+	return thread->ptid;
+    }
+
+  for (thread_info &thread : current_inferior ()->non_exited_threads ())
+    if (is_linux_thread (thread))
+      return thread.ptid;
 
   error (_("Could not find a non-exited thread to read /proc from."));
 }
diff --git a/gdb/testsuite/gdb.rocm/gcore-host.exp b/gdb/testsuite/gdb.rocm/gcore-host.exp
index 80bf60e4e9f5..867f803f4381 100644
--- a/gdb/testsuite/gdb.rocm/gcore-host.exp
+++ b/gdb/testsuite/gdb.rocm/gcore-host.exp
@@ -23,7 +23,10 @@ standard_testfile .cpp
 
 require allow_hip_tests
 
-if {[build_executable "failed to prepare" $testfile $srcfile {debug hip}]} {
+# Build a position-independent executable.  Loading the core file of a PIE
+# requires the auxiliary vector (NT_AUXV note) to find where the executable
+# is loaded.
+if {[build_executable "failed to prepare" $testfile $srcfile {debug hip pie}]} {
     return
 }
 
@@ -34,6 +37,15 @@ if {[build_executable "failed to prepare" $testfile $srcfile {debug hip}]} {
 proc do_test { selected_thread } {
     clean_restart $::testfile
 
+    # To verify that the auxiliary vector was properly written to the core,
+    # pass a long argument to the program, to make the command line longer than
+    # what the NT_PRPSINFO note can hold (80 characters).  When loading the
+    # core file, GDB uses the auxiliary vector to read the full command line
+    # from memory.  Without the auxiliary vector, it falls back to the
+    # truncated command line from the NT_PRPSINFO note.
+    set long_arg [string repeat x 100]
+    gdb_test_no_output "set args $long_arg"
+
     set corefile [standard_output_file $::testfile.$selected_thread.core]
 
     with_rocm_gpu_lock {
@@ -75,10 +87,21 @@ proc do_test { selected_thread } {
 
     clean_restart $::testfile
 
+    # Check that GDB shows the full command line, which requires the NT_AUXV
+    # note to be present and correct.
     gdb_test "core-file $corefile" \
-	"Core was generated by `[string_to_regexp $::binfile]'\\..*" \
+	"Core was generated by `[string_to_regexp "$::binfile $long_arg"]'\\..*" \
 	"load corefile"
 
+    # Check that the core file contains the auxiliary vector (NT_AUXV note).
+    gdb_test "info auxv" \
+	"AT_ENTRY +Entry point of program +$::hex\r\n.*AT_NULL +End of vector +0x0" \
+	"core file auxv"
+
+    # Check that GDB finds the shared libraries loaded by the program.
+    gdb_test "info sharedlibrary" "/libamdhip64\\.so.*" \
+	"core file shared libraries"
+
     # The core file is expected to contain all the host threads, and nothing
     # else.
     set num_core_threads 0
@@ -97,6 +120,6 @@ proc do_test { selected_thread } {
     }
 }
 
-foreach_with_prefix selected_thread { host } {
+foreach_with_prefix selected_thread { host wave } {
     do_test $selected_thread
 }
-- 
2.56.0