[PATCH 7/7] gdb/amdgpu: add follow fork and exec support

[Simon Marchi <simon.marchi@efficios.com>]

Prior to this patch, it's not possible for GDB to debug GPU code in fork
children or after an exec.  The amd-dbgapi target attaches to processes
when an inferior appears due to a "run" or "attach" command, but not
after a fork or exec.  This patch adds support for that, such that it's
possible to for an inferior to fork and for GDB to debug the GPU code in
the child.

To achieve that, use the inferior_forked and inferior_execd observers.

In the case of fork, we have nothing to do if `child_inf` is nullptr,
meaning that GDB won't debug the child.  We also don't attach if the
inferior has vforked.  We are already attached to the parent's address
space, which is shared with the child, so trying to attach would cause
problems.  And anyway, the inferior can't do anything other than exec or
exit, it certainly won't start GPU kernels before exec'ing.

In the case of exec, we detach from the exec'ing inferior and attach to
the following inferior.  This works regardless of whether they are the
same or not.  If they are the same, meaning the execution continues in
the existing inferior, we need to do a detach/attach anyway, as
amd-dbgapi needs to be aware of the new address space created by the
exec.

Note that we use observers and not target_ops::follow_{fork,exec} here.
When the amd-dbgapi target is compiled in, it will attach (in the
amd_dbgapi_process_attach sense, not the ptrace sense) to native
inferiors when they appear, but won't push itself on the inferior's
target stack just yet.  It only pushes itself if the inferior
initializes the ROCm runtime.  So, if a non-GPU-using inferior calls
fork, an amd_dbgapi_target::follow_fork method would not get called.
Same for exec.  A previous version of the code had the amd-dbgapi target
pushed all the time, in which case we could use the target methods.  But
we prefer having the target pushed only when necessary, it's less
intrusive when doing native debugging that doesn't involve the GPU.

Change-Id: I5819c151c371120da8bab2fa9cbfa8769ba1d6f9
Reviewed-By: Pedro Alves <pedro@palves.net>
---
 gdb/amd-dbgapi-target.c                       | 39 ++++++++
 .../fork-exec-gpu-to-non-gpu-execee.cpp       | 27 ++++++
 .../fork-exec-gpu-to-non-gpu-execer.cpp       | 55 ++++++++++++
 .../gdb.rocm/fork-exec-gpu-to-non-gpu.exp     | 89 +++++++++++++++++++
 .../fork-exec-non-gpu-to-gpu-execee.cpp       | 36 ++++++++
 .../fork-exec-non-gpu-to-gpu-execer.cpp       | 46 ++++++++++
 .../gdb.rocm/fork-exec-non-gpu-to-gpu.exp     | 88 ++++++++++++++++++
 7 files changed, 380 insertions(+)
 create mode 100644 gdb/testsuite/gdb.rocm/fork-exec-gpu-to-non-gpu-execee.cpp
 create mode 100644 gdb/testsuite/gdb.rocm/fork-exec-gpu-to-non-gpu-execer.cpp
 create mode 100644 gdb/testsuite/gdb.rocm/fork-exec-gpu-to-non-gpu.exp
 create mode 100644 gdb/testsuite/gdb.rocm/fork-exec-non-gpu-to-gpu-execee.cpp
 create mode 100644 gdb/testsuite/gdb.rocm/fork-exec-non-gpu-to-gpu-execer.cpp
 create mode 100644 gdb/testsuite/gdb.rocm/fork-exec-non-gpu-to-gpu.exp

diff --git a/gdb/amd-dbgapi-target.c b/gdb/amd-dbgapi-target.c
index f5161038c51d..61b9805abca0 100644
--- a/gdb/amd-dbgapi-target.c
+++ b/gdb/amd-dbgapi-target.c
@@ -1343,6 +1343,17 @@ attach_amd_dbgapi (inferior *inf)
       return;
     }
 
+  /* dbgapi can't attach to a vfork child (a process born from a vfork that
+     hasn't exec'ed yet) while we are still attached to the parent.  It would
+     not be useful for us to attach to vfork children anyway, because vfork
+     children are very restricted in what they can do (see vfork(2)) and aren't
+     going to launch some GPU programs that we need to debug.  To avoid this
+     problem, we don't push the amd-dbgapi target / attach dbgapi in vfork
+     children.  If a vfork child execs, we'll try enabling the amd-dbgapi target
+     through the inferior_execd observer.  */
+  if (inf->vfork_parent != nullptr)
+    return;
+
   auto *info = get_amd_dbgapi_inferior_info (inf);
 
   /* Are we already attached?  */
@@ -1655,6 +1666,32 @@ amd_dbgapi_target_inferior_created (inferior *inf)
   attach_amd_dbgapi (inf);
 }
 
+/* inferior_execd observer.  */
+
+static void
+amd_dbgapi_inferior_execd (inferior *exec_inf, inferior *follow_inf)
+{
+  /* The inferior has EXEC'd and the process image has changed.  The dbgapi is
+     attached to the old process image, so we need to detach and re-attach to
+     the new process image.  */
+  detach_amd_dbgapi (exec_inf);
+  attach_amd_dbgapi (follow_inf);
+}
+
+/* inferior_forked observer.  */
+
+static void
+amd_dbgapi_inferior_forked (inferior *parent_inf, inferior *child_inf,
+			    target_waitkind fork_kind)
+{
+  if (child_inf != nullptr  && fork_kind != TARGET_WAITKIND_VFORKED)
+    {
+      scoped_restore_current_thread restore_thread;
+      switch_to_thread (*child_inf->threads ().begin ());
+      attach_amd_dbgapi (child_inf);
+    }
+}
+
 /* inferior_exit observer.
 
    This covers normal exits, but also detached inferiors (including detached
@@ -1924,6 +1961,8 @@ _initialize_amd_dbgapi_target ()
   gdb::observers::inferior_created.attach
     (amd_dbgapi_target_inferior_created,
      amd_dbgapi_target_inferior_created_observer_token, "amd-dbgapi");
+  gdb::observers::inferior_execd.attach (amd_dbgapi_inferior_execd, "amd-dbgapi");
+  gdb::observers::inferior_forked.attach (amd_dbgapi_inferior_forked, "amd-dbgapi");
   gdb::observers::inferior_exit.attach (amd_dbgapi_inferior_exited, "amd-dbgapi");
   gdb::observers::inferior_pre_detach.attach (amd_dbgapi_inferior_pre_detach, "amd-dbgapi");
 
diff --git a/gdb/testsuite/gdb.rocm/fork-exec-gpu-to-non-gpu-execee.cpp b/gdb/testsuite/gdb.rocm/fork-exec-gpu-to-non-gpu-execee.cpp
new file mode 100644
index 000000000000..eacfcd86faa8
--- /dev/null
+++ b/gdb/testsuite/gdb.rocm/fork-exec-gpu-to-non-gpu-execee.cpp
@@ -0,0 +1,27 @@
+/* This testcase is part of GDB, the GNU debugger.
+
+   Copyright 2021-2023 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/>.  */
+
+static void
+break_here_execee (void)
+{}
+
+int
+main (void)
+{
+  break_here_execee ();
+  return 0;
+}
diff --git a/gdb/testsuite/gdb.rocm/fork-exec-gpu-to-non-gpu-execer.cpp b/gdb/testsuite/gdb.rocm/fork-exec-gpu-to-non-gpu-execer.cpp
new file mode 100644
index 000000000000..1a731aeca517
--- /dev/null
+++ b/gdb/testsuite/gdb.rocm/fork-exec-gpu-to-non-gpu-execer.cpp
@@ -0,0 +1,55 @@
+/* This testcase is part of GDB, the GNU debugger.
+
+   Copyright 2021-2023 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 <unistd.h>
+
+__global__ static void
+kernel1 ()
+{}
+
+__device__ static void
+break_here_execer ()
+{
+}
+
+__global__ static void
+kernel2 ()
+{
+  break_here_execer ();
+}
+
+int
+main ()
+{
+  /* Launch a first kernel to make sure the runtime is active by the time we
+     call fork.  */
+  kernel1<<<1, 1>>> ();
+
+  /* fork + exec while the runtime is active.  */
+  if (FORK () == 0)
+    {
+      int ret = execl (EXECEE, EXECEE, NULL);
+      perror ("exec");
+      abort ();
+    }
+
+  kernel2<<<1, 1>>> ();
+
+  hipDeviceSynchronize ();
+  return 0;
+}
diff --git a/gdb/testsuite/gdb.rocm/fork-exec-gpu-to-non-gpu.exp b/gdb/testsuite/gdb.rocm/fork-exec-gpu-to-non-gpu.exp
new file mode 100644
index 000000000000..852294b7067b
--- /dev/null
+++ b/gdb/testsuite/gdb.rocm/fork-exec-gpu-to-non-gpu.exp
@@ -0,0 +1,89 @@
+# Copyright 2021-2023 Free Software Foundation, Inc.
+
+# 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/>.
+
+# Verify handling of a GPU program that does a (v)fork + exec to execute
+# a non-GPU program.
+
+load_lib rocm.exp
+
+require allow_hipcc_tests
+
+standard_testfile -execer.cpp -execee.cpp
+
+set srcfile_execer "$srcfile"
+set srcfile_execee "$srcfile2"
+set binfile_execee "$binfile-execee"
+
+# Compile two versions of execer, one that uses fork and one that uses vfork.
+foreach_with_prefix fork_func { fork vfork } {
+    set opts [list debug hip additional_flags=-DFORK=$fork_func \
+	additional_flags=-DEXECEE="${::binfile_execee}"]
+    if {[build_executable "failed to prepare" ${::binfile}-execer-${fork_func} \
+	    $srcfile_execer $opts]} {
+	return
+    }
+}
+
+if {[build_executable "failed to prepare" $binfile_execee $srcfile_execee \
+	{debug}]} {
+    return
+}
+
+proc do_test { detach-on-fork follow-fork-mode fork_func } {
+    # In this case, the parent can't execute, as it's blocked in
+    # vfork.  Skip it.
+    if { ${detach-on-fork} == "off"
+	 && ${follow-fork-mode} == "parent"
+	 && ${fork_func} == "vfork" } {
+	return
+    }
+
+    with_rocm_gpu_lock {
+	clean_restart ${::binfile}-execer-${fork_func}
+
+	gdb_test_no_output "set detach-on-fork ${detach-on-fork}"
+	gdb_test_no_output "set follow-fork-mode ${follow-fork-mode}"
+
+	if { ${follow-fork-mode} == "parent" } {
+	    runto break_here_execer allow-pending message
+	    gdb_continue_to_end "continue parent to end" "continue" 1
+
+	    if { ${detach-on-fork} == "off" } {
+		gdb_test "inferior 2" "Switching to inferior 2 .*"
+		gdb_continue_to_end "continue child to end" "continue" 1
+	    }
+	} elseif { ${follow-fork-mode} == "child" } {
+	    runto break_here_execee allow-pending message
+	    gdb_continue_to_end "continue child to end" "continue" 1
+
+	    if { ${detach-on-fork} == "off" } {
+		gdb_test "inferior 1" "Switching to inferior 1 .*"
+		gdb_continue_to_end "continue parent to end" "continue" 1
+	    }
+	} else {
+	    error "unexpected follow-fork-mode value: ${follow-fork-mode}"
+	}
+    }
+}
+
+foreach_with_prefix detach-on-fork { on off } {
+    foreach_with_prefix follow-fork-mode { parent child } {
+	foreach_with_prefix fork_func { fork vfork } {
+	    do_test ${detach-on-fork} ${follow-fork-mode} $fork_func
+	}
+    }
+}
diff --git a/gdb/testsuite/gdb.rocm/fork-exec-non-gpu-to-gpu-execee.cpp b/gdb/testsuite/gdb.rocm/fork-exec-non-gpu-to-gpu-execee.cpp
new file mode 100644
index 000000000000..2de8fe20a0d6
--- /dev/null
+++ b/gdb/testsuite/gdb.rocm/fork-exec-non-gpu-to-gpu-execee.cpp
@@ -0,0 +1,36 @@
+/* This testcase is part of GDB, the GNU debugger.
+
+   Copyright 2021-2023 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>
+
+__device__ static void
+break_here_execee ()
+{}
+
+__global__ void
+kernel ()
+{
+  break_here_execee ();
+}
+
+int
+main ()
+{
+  kernel<<<1, 1>>> ();
+  hipDeviceSynchronize ();
+  return 0;
+}
diff --git a/gdb/testsuite/gdb.rocm/fork-exec-non-gpu-to-gpu-execer.cpp b/gdb/testsuite/gdb.rocm/fork-exec-non-gpu-to-gpu-execer.cpp
new file mode 100644
index 000000000000..3ee07949273a
--- /dev/null
+++ b/gdb/testsuite/gdb.rocm/fork-exec-non-gpu-to-gpu-execer.cpp
@@ -0,0 +1,46 @@
+/* This testcase is part of GDB, the GNU debugger.
+
+   Copyright 2021-2023 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 <sys/types.h>
+#include <unistd.h>
+#include <stdlib.h>
+#include <stdio.h>
+
+static void
+break_here_execer ()
+{}
+
+int
+main ()
+{
+  /* FORK is defined to fork or vfork by the test.  */
+  int pid = FORK ();
+  if (pid != 0)
+    {
+      /* Parent.  */
+      break_here_execer ();
+    }
+  else
+    {
+      /* EXECEE is defined by the test.  */
+      int ret = execl (EXECEE, EXECEE, NULL);
+      perror ("exec");
+      abort ();
+    }
+
+    return 0;
+}
diff --git a/gdb/testsuite/gdb.rocm/fork-exec-non-gpu-to-gpu.exp b/gdb/testsuite/gdb.rocm/fork-exec-non-gpu-to-gpu.exp
new file mode 100644
index 000000000000..e372db5a32e6
--- /dev/null
+++ b/gdb/testsuite/gdb.rocm/fork-exec-non-gpu-to-gpu.exp
@@ -0,0 +1,88 @@
+# Copyright 2021-2023 Free Software Foundation, Inc.
+
+# 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/>.
+
+# Verify that we can debug a GPU program in a child after a (v)fork + exec.
+
+load_lib rocm.exp
+
+require allow_hipcc_tests
+
+standard_testfile -execer.cpp -execee.cpp
+
+set srcfile_execer "$srcfile"
+set srcfile_execee "$srcfile2"
+set binfile_execee "$binfile-execee"
+
+# Compile two versions of execer, one that uses fork and one that uses vfork.
+foreach_with_prefix fork_func { fork vfork } {
+    set opts [list additional_flags=-DFORK=$fork_func \
+	additional_flags=-DEXECEE="${::binfile_execee}"]
+    if {[build_executable "failed to prepare" ${::binfile}-execer-${fork_func} \
+	    $srcfile_execer $opts]} {
+	return
+    }
+}
+
+if {[build_executable "failed to prepare" $binfile_execee $srcfile_execee \
+	{debug hip}]} {
+    return
+}
+
+proc do_test { detach-on-fork follow-fork-mode fork_func } {
+    # In this case, the parent can't execute, as it's blocked in
+    # vfork.  Skip it.
+    if { ${detach-on-fork} == "off"
+	 && ${follow-fork-mode} == "parent"
+	 && ${fork_func} == "vfork" } {
+	return
+    }
+
+    with_rocm_gpu_lock {
+	clean_restart ${::binfile}-execer-${fork_func}
+
+	gdb_test_no_output "set detach-on-fork ${detach-on-fork}"
+	gdb_test_no_output "set follow-fork-mode ${follow-fork-mode}"
+
+	if { ${follow-fork-mode} == "parent" } {
+	    runto break_here_execer allow-pending message
+	    gdb_continue_to_end "continue parent to end" "continue" 1
+
+	    if { ${detach-on-fork} == "off" } {
+		gdb_test "inferior 2" "Switching to inferior 2 .*"
+		gdb_continue_to_end "continue child to end" "continue" 1
+	    }
+	} elseif { ${follow-fork-mode} == "child" } {
+	    runto break_here_execee allow-pending message
+	    gdb_continue_to_end "continue child to end" "continue" 1
+
+	    if { ${detach-on-fork} == "off" } {
+		gdb_test "inferior 1" "Switching to inferior 1 .*"
+		gdb_continue_to_end "continue parent to end" "continue" 1
+	    }
+	} else {
+	    error "unexpected follow-fork-mode value: ${follow-fork-mode}"
+	}
+    }
+}
+
+foreach_with_prefix detach-on-fork { on off } {
+    foreach_with_prefix follow-fork-mode { parent child } {
+	foreach_with_prefix fork_func { fork vfork } {
+	    do_test ${detach-on-fork} ${follow-fork-mode} $fork_func
+	}
+    }
+}
-- 
2.40.0