[PATCH v2 1/6] gdb/amd-dbgapi: remove unnecessary AMD_DBGAPI_EVENT_KIND_NONE argument

[PATCH v2 1/6] gdb/amd-dbgapi: remove unnecessary AMD_DBGAPI_EVENT_KIND_NONE argument

[Simon Marchi]

Rely on the default value.

Change-Id: I08c683de005806c5c5d29ed7f9b0c6de81b49a01
Approved-By: Lancelot Six <lancelot.six@amd.com> (amdgpu)
---
 gdb/amd-dbgapi-target.c | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/gdb/amd-dbgapi-target.c b/gdb/amd-dbgapi-target.c
index 819b373da321..2aa9d2a03a4b 100644
--- a/gdb/amd-dbgapi-target.c
+++ b/gdb/amd-dbgapi-target.c
@@ -1473,7 +1473,7 @@ amd_dbgapi_target::wait (ptid_t ptid, struct target_waitstatus *ws,
       /* Drain the events for the current inferior from the amd_dbgapi and
 	 preserve the ordering.  */
       auto info = get_amd_dbgapi_inferior_info (current_inferior ());
-      process_event_queue (info->process_id, AMD_DBGAPI_EVENT_KIND_NONE);
+      process_event_queue (info->process_id);
 
       std::tie (event_ptid, gpu_waitstatus) = consume_one_event (ptid.pid ());
       if (event_ptid == minus_one_ptid)

base-commit: ab65e51fa9d3b9ca44944ce191297ee4f23a33e6
-- 
2.49.0

[PATCH v2 2/6] gdb/amd-dbgapi: add assert in require_forward_progress

[Simon Marchi]

I didn't have a problem in this area, but it seems to me that this
pre-condition should always hold.  We should only disable forward
progress requirement if the target says it's ok to do so.  Otherwise, we
could get in a situation where we wait for events from amd-dbgapi, which
will never arrive, because amd-dbgapi didn't actually resume things.

Change-Id: Ifc49f55c7874924b7c47888b8391a07a01d960fc
Approved-by: Lancelot Six <lancelot.six@amd.com>
---
 gdb/amd-dbgapi-target.c | 6 ++++++
 1 file changed, 6 insertions(+)

diff --git a/gdb/amd-dbgapi-target.c b/gdb/amd-dbgapi-target.c
index 2aa9d2a03a4b..32a3f6d3c334 100644
--- a/gdb/amd-dbgapi-target.c
+++ b/gdb/amd-dbgapi-target.c
@@ -450,6 +450,12 @@ static void
 require_forward_progress (ptid_t ptid, process_stratum_target *proc_target,
 			  bool require)
 {
+  /* If we try to disable forward progress requirement but the target expects
+     resumed threads to be committed to the target, we could wait for events
+     that will never arrive.  */
+  if (!require)
+    gdb_assert (!proc_target->commit_resumed_state);
+
   for (inferior *inf : all_inferiors (proc_target))
     {
       if (ptid != minus_one_ptid && inf->pid != ptid.pid ())
-- 
2.49.0

[PATCH v2 3/6] gdb/amd-dbgapi: pass amd_dbgapi_inferior_info to process_event_queue

[Simon Marchi]

New in v2: update a comment

A following patch will make process_event_queue access a field of
amd_dbgapi_inferior_info.  Prepare for this by making
process_event_queue accept an amd_dbgapi_inferior_info object, instead
of a process id.

Change-Id: I9adc491dd1ff64ff74c40aa7662fffb11bd8332b
---
 gdb/amd-dbgapi-target.c | 25 +++++++++----------------
 1 file changed, 9 insertions(+), 16 deletions(-)

diff --git a/gdb/amd-dbgapi-target.c b/gdb/amd-dbgapi-target.c
index 32a3f6d3c334..888552e00d3e 100644
--- a/gdb/amd-dbgapi-target.c
+++ b/gdb/amd-dbgapi-target.c
@@ -234,7 +234,7 @@ struct amd_dbgapi_inferior_info
 };
 
 static amd_dbgapi_event_id_t process_event_queue
-  (amd_dbgapi_process_id_t process_id,
+  (amd_dbgapi_inferior_info &info,
    amd_dbgapi_event_kind_t until_event_kind = AMD_DBGAPI_EVENT_KIND_NONE);
 
 static const target_info amd_dbgapi_target_info = {
@@ -564,8 +564,7 @@ amd_dbgapi_target_breakpoint::check_status (struct bpstat *bs)
   /* If the action is AMD_DBGAPI_BREAKPOINT_ACTION_HALT, we need to wait until
      a breakpoint resume event for this breakpoint_id is seen.  */
   amd_dbgapi_event_id_t resume_event_id
-    = process_event_queue (info->process_id,
-			   AMD_DBGAPI_EVENT_KIND_BREAKPOINT_RESUME);
+    = process_event_queue (*info, AMD_DBGAPI_EVENT_KIND_BREAKPOINT_RESUME);
 
   /* We should always get a breakpoint_resume event after processing all
      events generated by reporting the breakpoint hit.  */
@@ -1337,28 +1336,22 @@ event_kind_str (amd_dbgapi_event_kind_t kind)
   gdb_assert_not_reached ("unhandled amd_dbgapi_event_kind_t value");
 }
 
-/* Drain the dbgapi event queue of a given process_id, or of all processes if
-   process_id is AMD_DBGAPI_PROCESS_NONE.  Stop processing the events if an
-   event of a given kind is requested and `process_id` is not
-   AMD_DBGAPI_PROCESS_NONE.  Wave stop events that are not returned are queued
-   into their inferior's amd_dbgapi_inferior_info pending wave events. */
+/* Drain the dbgapi event queue of a given inferior.  Stop processing the
+   events if an event of a given kind is requested (not AMD_DBGAPI_EVENT_NONE).
+   Wave stop events that are not returned are queued into their inferior's
+   amd_dbgapi_inferior_info pending wave events. */
 
 static amd_dbgapi_event_id_t
-process_event_queue (amd_dbgapi_process_id_t process_id,
+process_event_queue (amd_dbgapi_inferior_info &info,
 		     amd_dbgapi_event_kind_t until_event_kind)
 {
-  /* An event of a given type can only be requested from a single
-     process_id.  */
-  gdb_assert (until_event_kind == AMD_DBGAPI_EVENT_KIND_NONE
-	      || process_id != AMD_DBGAPI_PROCESS_NONE);
-
   while (true)
     {
       amd_dbgapi_event_id_t event_id;
       amd_dbgapi_event_kind_t event_kind;
 
       amd_dbgapi_status_t status
-	= amd_dbgapi_process_next_pending_event (process_id, &event_id,
+	= amd_dbgapi_process_next_pending_event (info.process_id, &event_id,
 						 &event_kind);
 
       if (status != AMD_DBGAPI_STATUS_SUCCESS)
@@ -1479,7 +1472,7 @@ amd_dbgapi_target::wait (ptid_t ptid, struct target_waitstatus *ws,
       /* Drain the events for the current inferior from the amd_dbgapi and
 	 preserve the ordering.  */
       auto info = get_amd_dbgapi_inferior_info (current_inferior ());
-      process_event_queue (info->process_id);
+      process_event_queue (*info);
 
       std::tie (event_ptid, gpu_waitstatus) = consume_one_event (ptid.pid ());
       if (event_ptid == minus_one_ptid)
-- 
2.49.0

[PATCH v2 4/6] gdb/amd-dbgapi: pass amd_dbgapi_inferior_info to process_one_event

[Simon Marchi]

New in v2: this patch is new

Pass the amd_dbgapi_inferior_info object from process_event_queue to
process_one_event.  Since process_event_queue pulls events for one
specific inferior, we know for which inferior the event is.  This
removes the need for process_one_event to do two dbgapi calls to get the
relevant pid.  If also removes one inferior lookup.

Change-Id: I22927e4b6251513eb3be95785082058aa3d09954
---
 gdb/amd-dbgapi-target.c | 58 +++++++++++++++--------------------------
 1 file changed, 21 insertions(+), 37 deletions(-)

diff --git a/gdb/amd-dbgapi-target.c b/gdb/amd-dbgapi-target.c
index 888552e00d3e..80111359c8cd 100644
--- a/gdb/amd-dbgapi-target.c
+++ b/gdb/amd-dbgapi-target.c
@@ -1143,32 +1143,14 @@ add_gpu_thread (inferior *inf, ptid_t wave_ptid)
 /* Process an event that was just pulled out of the amd-dbgapi library.  */
 
 static void
-process_one_event (amd_dbgapi_event_id_t event_id,
+process_one_event (amd_dbgapi_inferior_info &info,
+		   amd_dbgapi_event_id_t event_id,
 		   amd_dbgapi_event_kind_t event_kind)
 {
   /* Automatically mark this event processed when going out of scope.  */
   scoped_amd_dbgapi_event_processed mark_event_processed (event_id);
 
-  amd_dbgapi_process_id_t process_id;
-  amd_dbgapi_status_t status
-    = amd_dbgapi_event_get_info (event_id, AMD_DBGAPI_EVENT_INFO_PROCESS,
-				 sizeof (process_id), &process_id);
-  if (status != AMD_DBGAPI_STATUS_SUCCESS)
-    error (_("event_get_info for event_%ld failed (%s)"), event_id.handle,
-	   get_status_string (status));
-
-  amd_dbgapi_os_process_id_t pid;
-  status = amd_dbgapi_process_get_info (process_id,
-					AMD_DBGAPI_PROCESS_INFO_OS_ID,
-					sizeof (pid), &pid);
-  if (status != AMD_DBGAPI_STATUS_SUCCESS)
-    error (_("process_get_info for process_%ld failed (%s)"),
-	   process_id.handle, get_status_string (status));
-
-  auto *proc_target = current_inferior ()->process_target ();
-  inferior *inf = find_inferior_pid (proc_target, pid);
-  gdb_assert (inf != nullptr);
-  amd_dbgapi_inferior_info *info = get_amd_dbgapi_inferior_info (inf);
+  gdb_assert (info.inf != nullptr);
 
   switch (event_kind)
     {
@@ -1176,14 +1158,14 @@ process_one_event (amd_dbgapi_event_id_t event_id,
     case AMD_DBGAPI_EVENT_KIND_WAVE_STOP:
       {
 	amd_dbgapi_wave_id_t wave_id;
-	status
+	amd_dbgapi_status_t status
 	  = amd_dbgapi_event_get_info (event_id, AMD_DBGAPI_EVENT_INFO_WAVE,
 				       sizeof (wave_id), &wave_id);
 	if (status != AMD_DBGAPI_STATUS_SUCCESS)
 	  error (_("event_get_info for event_%ld failed (%s)"),
 		 event_id.handle, get_status_string (status));
 
-	ptid_t event_ptid = make_gpu_ptid (pid, wave_id);
+	ptid_t event_ptid = make_gpu_ptid (info.inf->pid, wave_id);
 	target_waitstatus ws;
 
 	amd_dbgapi_wave_stop_reasons_t stop_reason;
@@ -1224,9 +1206,10 @@ process_one_event (amd_dbgapi_event_id_t event_id,
 	    else
 	      ws.set_stopped (GDB_SIGNAL_0);
 
-	    thread_info *thread = proc_target->find_thread (event_ptid);
-	    if (thread == nullptr)
-	      thread = add_gpu_thread (inf, event_ptid);
+            thread_info *thread =
+                info.inf->process_target()->find_thread(event_ptid);
+            if (thread == nullptr)
+	      thread = add_gpu_thread (info.inf, event_ptid);
 
 	    /* If the wave is stopped because of a software breakpoint, the
 	       program counter needs to be adjusted so that it points to the
@@ -1248,7 +1231,7 @@ process_one_event (amd_dbgapi_event_id_t event_id,
 	  error (_("wave_get_info for wave_%ld failed (%s)"),
 		 wave_id.handle, get_status_string (status));
 
-	info->wave_events.emplace_back (event_ptid, ws);
+	info.wave_events.emplace_back (event_ptid, ws);
 	break;
       }
 
@@ -1266,7 +1249,7 @@ process_one_event (amd_dbgapi_event_id_t event_id,
 	 When amd_dbgapi_target_breakpoint::check_status is called, the current
 	 inferior is the inferior that hit the breakpoint, which should still be
 	 the case now.  */
-      gdb_assert (inf == current_inferior ());
+      gdb_assert (info.inf == current_inferior ());
       handle_solib_event ();
       break;
 
@@ -1280,22 +1263,23 @@ process_one_event (amd_dbgapi_event_id_t event_id,
       {
 	amd_dbgapi_runtime_state_t runtime_state;
 
-	status = amd_dbgapi_event_get_info (event_id,
-					    AMD_DBGAPI_EVENT_INFO_RUNTIME_STATE,
-					    sizeof (runtime_state),
-					    &runtime_state);
+	amd_dbgapi_status_t status
+	  = amd_dbgapi_event_get_info (event_id,
+				       AMD_DBGAPI_EVENT_INFO_RUNTIME_STATE,
+				       sizeof (runtime_state),
+				       &runtime_state);
 	if (status != AMD_DBGAPI_STATUS_SUCCESS)
 	  error (_("event_get_info for event_%ld failed (%s)"),
 		 event_id.handle, get_status_string (status));
 
 	gdb_assert (runtime_state == AMD_DBGAPI_RUNTIME_STATE_UNLOADED);
 	gdb_assert
-	  (info->runtime_state == AMD_DBGAPI_RUNTIME_STATE_LOADED_SUCCESS);
+	  (info.runtime_state == AMD_DBGAPI_RUNTIME_STATE_LOADED_SUCCESS);
 
-	info->runtime_state = runtime_state;
+	info.runtime_state = runtime_state;
 
-	gdb_assert (inf->target_is_pushed (&the_amd_dbgapi_target));
-	inf->unpush_target (&the_amd_dbgapi_target);
+	gdb_assert (info.inf->target_is_pushed (&the_amd_dbgapi_target));
+	info.inf->unpush_target (&the_amd_dbgapi_target);
       }
       break;
 
@@ -1367,7 +1351,7 @@ process_event_queue (amd_dbgapi_inferior_info &info,
       if (event_id == AMD_DBGAPI_EVENT_NONE || event_kind == until_event_kind)
 	return event_id;
 
-      process_one_event (event_id, event_kind);
+      process_one_event (info, event_id, event_kind);
     }
 }
 
-- 
2.49.0

[PATCH v2 5/6] gdb/amd-dbgapi: factor out require_forward_progress overload to target one inferior

[Simon Marchi]

A following patch will want to call require_forward_progress for a given
inferior.  Extract a new require_forward_progress overload from the
existing require_forward_progress function that targets a specific
inferior.

Change-Id: I54f42b83eb8443d4d91747ffbc86eaeb017f1e49
Approved-by: Lancelot Six <lancelot.six@amd.com>
---
 gdb/amd-dbgapi-target.c | 49 +++++++++++++++++++++++------------------
 1 file changed, 28 insertions(+), 21 deletions(-)

diff --git a/gdb/amd-dbgapi-target.c b/gdb/amd-dbgapi-target.c
index 80111359c8cd..2512083db5d9 100644
--- a/gdb/amd-dbgapi-target.c
+++ b/gdb/amd-dbgapi-target.c
@@ -443,6 +443,32 @@ async_event_handler_mark ()
   mark_async_event_handler (amd_dbgapi_async_event_handler);
 }
 
+/* Set forward progress requirement to REQUIRE for inferior INFO.  */
+
+static void
+require_forward_progress (amd_dbgapi_inferior_info &info, bool require)
+{
+  /* If we try to disable forward progress requirement but the target expects
+     resumed threads to be committed to the target, we could wait for events
+     that will never arrive.  */
+  if (!require)
+    gdb_assert (!info.inf->process_target ()->commit_resumed_state);
+
+  gdb_assert (info.process_id != AMD_DBGAPI_PROCESS_NONE);
+
+  /* Don't do unnecessary calls to amd-dbgapi to avoid polluting the logs.  */
+  if (info.forward_progress_required == require)
+    return;
+
+  const auto progress
+    = require ? AMD_DBGAPI_PROGRESS_NORMAL : AMD_DBGAPI_PROGRESS_NO_FORWARD;
+  const auto status
+    = amd_dbgapi_process_set_progress (info.process_id, progress);
+  gdb_assert (status == AMD_DBGAPI_STATUS_SUCCESS);
+
+  info.forward_progress_required = require;
+}
+
 /* Set forward progress requirement to REQUIRE for all processes of PROC_TARGET
    matching PTID.  */
 
@@ -450,12 +476,6 @@ static void
 require_forward_progress (ptid_t ptid, process_stratum_target *proc_target,
 			  bool require)
 {
-  /* If we try to disable forward progress requirement but the target expects
-     resumed threads to be committed to the target, we could wait for events
-     that will never arrive.  */
-  if (!require)
-    gdb_assert (!proc_target->commit_resumed_state);
-
   for (inferior *inf : all_inferiors (proc_target))
     {
       if (ptid != minus_one_ptid && inf->pid != ptid.pid ())
@@ -463,21 +483,8 @@ require_forward_progress (ptid_t ptid, process_stratum_target *proc_target,
 
       amd_dbgapi_inferior_info *info = get_amd_dbgapi_inferior_info (inf);
 
-      if (info->process_id == AMD_DBGAPI_PROCESS_NONE)
-	continue;
-
-      /* Don't do unnecessary calls to amd-dbgapi to avoid polluting the logs.  */
-      if (info->forward_progress_required == require)
-	continue;
-
-      amd_dbgapi_status_t status
-	= amd_dbgapi_process_set_progress
-	    (info->process_id, (require
-				? AMD_DBGAPI_PROGRESS_NORMAL
-				: AMD_DBGAPI_PROGRESS_NO_FORWARD));
-      gdb_assert (status == AMD_DBGAPI_STATUS_SUCCESS);
-
-      info->forward_progress_required = require;
+      if (info->process_id != AMD_DBGAPI_PROCESS_NONE)
+	require_forward_progress (*info, require);
 
       /* If ptid targets a single inferior and we have found it, no need to
 	 continue.  */
-- 
2.49.0

[PATCH v2 6/6] gdb/amd-dbgapi: disable forward progress requirement in amd_dbgapi_target_breakpoint::check_status

[Simon Marchi]

New in v2:

  - fix / add copyright notice
  - remove unnecessary asm intruction in the test

ROCgdb handles target events very slowly when running a test case like
this, where a breakpoint is preset on HipTest::vectorADD:

    for (int i=0; i < numDevices; ++i) {
      HIPCHECK(hipSetDevice(i));
      hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, stream[i],
                        static_cast<const int*>(A_d[i]), static_cast<const int*>(B_d[i]), C_d[i], N);
    }

What happens is:

 - A kernel is launched
 - The internal runtime breakpoint is hit during the second
   hipLaunchKernelGGL call, which causes
   amd_dbgapi_target_breakpoint::check_status to be called
 - Meanwhile, all waves of the kernel hit the breakpoint on vectorADD
 - amd_dbgapi_target_breakpoint::check_status calls process_event_queue,
   which pulls the thousand of breakpoint hit events from the kernel
 - As part of handling the breakpoint hit events, we write the PC of the
   waves that stopped to decrement it.  Because the forward progress
   requirement is not disabled, this causes a suspend/resume of the
   queue each time, which is time-consuming.

The stack trace where this all happens is:

    #32 0x00007ffff6b9abda in amd_dbgapi_write_register (wave_id=..., register_id=..., offset=0, value_size=8, value=0x7fffea9fdcc0) at /home/smarchi/src/amd-dbgapi/src/register.cpp:587
    #33 0x00005555588c0bed in amd_dbgapi_target::store_registers (this=0x55555c7b1d20 <the_amd_dbgapi_target>, regcache=0x507000002240, regno=470) at /home/smarchi/src/wt/amd/gdb/amd-dbgapi-target.c:2504
    #34 0x000055555a5186a1 in target_store_registers (regcache=0x507000002240, regno=470) at /home/smarchi/src/wt/amd/gdb/target.c:3973
    #35 0x0000555559fab831 in regcache::raw_write (this=0x507000002240, regnum=470, src=...) at /home/smarchi/src/wt/amd/gdb/regcache.c:890
    #36 0x0000555559fabd2b in regcache::cooked_write (this=0x507000002240, regnum=470, src=...) at /home/smarchi/src/wt/amd/gdb/regcache.c:915
    #37 0x0000555559fc3ca5 in regcache::cooked_write<unsigned long, void> (this=0x507000002240, regnum=470, val=140737323456768) at /home/smarchi/src/wt/amd/gdb/regcache.c:850
    #38 0x0000555559fab09a in regcache_cooked_write_unsigned (regcache=0x507000002240, regnum=470, val=140737323456768) at /home/smarchi/src/wt/amd/gdb/regcache.c:858
    #39 0x0000555559fb0678 in regcache_write_pc (regcache=0x507000002240, pc=0x7ffff62bd900) at /home/smarchi/src/wt/amd/gdb/regcache.c:1460
    #40 0x00005555588bb37d in process_one_event (event_id=..., event_kind=AMD_DBGAPI_EVENT_KIND_WAVE_STOP) at /home/smarchi/src/wt/amd/gdb/amd-dbgapi-target.c:1873
    #41 0x00005555588bbf7b in process_event_queue (process_id=..., until_event_kind=AMD_DBGAPI_EVENT_KIND_BREAKPOINT_RESUME) at /home/smarchi/src/wt/amd/gdb/amd-dbgapi-target.c:2006
    #42 0x00005555588b1aca in amd_dbgapi_target_breakpoint::check_status (this=0x511000140900, bs=0x50600014ed00) at /home/smarchi/src/wt/amd/gdb/amd-dbgapi-target.c:890
    #43 0x0000555558c50080 in bpstat_stop_status (aspace=0x5070000061b0, bp_addr=0x7fffed0b9ab0, thread=0x518000026c80, ws=..., stop_chain=0x50600014ed00) at /home/smarchi/src/wt/amd/gdb/breakpoint.c:6126
    #44 0x000055555984f4ff in handle_signal_stop (ecs=0x7fffeaa40ef0) at /home/smarchi/src/wt/amd/gdb/infrun.c:7169
    #45 0x000055555984b889 in handle_inferior_event (ecs=0x7fffeaa40ef0) at /home/smarchi/src/wt/amd/gdb/infrun.c:6621
    #46 0x000055555983eab6 in fetch_inferior_event () at /home/smarchi/src/wt/amd/gdb/infrun.c:4750
    #47 0x00005555597caa5f in inferior_event_handler (event_type=INF_REG_EVENT) at /home/smarchi/src/wt/amd/gdb/inf-loop.c:42
    #48 0x00005555588b838e in handle_target_event (client_data=0x0) at /home/smarchi/src/wt/amd/gdb/amd-dbgapi-target.c:1513

Fix that performance problem by disabling the forward progress
requirement in amd_dbgapi_target_breakpoint::check_status, before
calling process_event_queue, so that we can process all events
efficiently.

Since the same performance problem could theoritically happen any time
process_event_queue is called with forward progress requirement enabled,
add an assert to ensure that forward progress requirement is disabled
when process_event_queue is invoked.  This makes it necessary to add a
require_forward_progress call to amd_dbgapi_finalize_core_attach.  It
looks a bit strange, since core files don't have execution, but it
doesn't hurt.

Add a test that replicates this scenario.  The test launches a kernel
that hits a breakpoint (with an always false condition) repeatedly.
Meanwhile, the host process loads an unloads a code object, causing
check_status to be called.

Bug: SWDEV-482511
Change-Id: Ida86340d679e6bd8462712953458c07ba3fd49ec
Approved-by: Lancelot Six <lancelot.six@amd.com>
---
 gdb/amd-dbgapi-target.c                       |  6 ++
 .../code-object-load-while-breakpoint-hit.cpp | 86 +++++++++++++++++++
 .../code-object-load-while-breakpoint-hit.exp | 68 +++++++++++++++
 3 files changed, 160 insertions(+)
 create mode 100644 gdb/testsuite/gdb.rocm/code-object-load-while-breakpoint-hit.cpp
 create mode 100644 gdb/testsuite/gdb.rocm/code-object-load-while-breakpoint-hit.exp

diff --git a/gdb/amd-dbgapi-target.c b/gdb/amd-dbgapi-target.c
index 2512083db5d9..4538cde87c46 100644
--- a/gdb/amd-dbgapi-target.c
+++ b/gdb/amd-dbgapi-target.c
@@ -568,6 +568,8 @@ amd_dbgapi_target_breakpoint::check_status (struct bpstat *bs)
   if (action == AMD_DBGAPI_BREAKPOINT_ACTION_RESUME)
     return;
 
+  require_forward_progress (*info, false);
+
   /* If the action is AMD_DBGAPI_BREAKPOINT_ACTION_HALT, we need to wait until
      a breakpoint resume event for this breakpoint_id is seen.  */
   amd_dbgapi_event_id_t resume_event_id
@@ -1336,6 +1338,10 @@ static amd_dbgapi_event_id_t
 process_event_queue (amd_dbgapi_inferior_info &info,
 		     amd_dbgapi_event_kind_t until_event_kind)
 {
+  /* Pulling events with forward progress required may result in bad
+     performance, make sure it is not required.  */
+  gdb_assert (!info.forward_progress_required);
+
   while (true)
     {
       amd_dbgapi_event_id_t event_id;
diff --git a/gdb/testsuite/gdb.rocm/code-object-load-while-breakpoint-hit.cpp b/gdb/testsuite/gdb.rocm/code-object-load-while-breakpoint-hit.cpp
new file mode 100644
index 000000000000..d75bc76fda4e
--- /dev/null
+++ b/gdb/testsuite/gdb.rocm/code-object-load-while-breakpoint-hit.cpp
@@ -0,0 +1,86 @@
+/* This testcase is part of GDB, the GNU debugger.
+
+   Copyright 2025 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/>.  */
+
+#ifdef DEVICE
+
+#include <hip/hip_runtime.h>
+
+constexpr unsigned int NUM_BREAKPOINT_HITS = 5;
+
+static __device__ void
+break_here ()
+{
+}
+
+extern "C" __global__ void
+kernel ()
+{
+  for (int n = 0; n < NUM_BREAKPOINT_HITS; ++n)
+    break_here ();
+}
+
+#else
+
+#include <hip/hip_runtime.h>
+#include <unistd.h>
+
+constexpr unsigned int NUM_ITEMS_PER_BLOCK = 256;
+constexpr unsigned int NUM_BLOCKS = 128;
+constexpr unsigned int NUM_ITEMS = NUM_ITEMS_PER_BLOCK * NUM_BLOCKS;
+constexpr unsigned int NUM_LOAD_UNLOADS = 5;
+
+#define CHECK(cmd)                                                            \
+  {                                                                           \
+    hipError_t error = cmd;                                                   \
+    if (error != hipSuccess)                                                  \
+      {                                                                       \
+	fprintf (stderr, "error: '%s'(%d) at %s:%d\n",                        \
+		 hipGetErrorString (error), error, __FILE__, __LINE__);       \
+	exit (EXIT_FAILURE);                                                  \
+      }                                                                       \
+  }
+
+int
+main (int argc, const char **argv)
+{
+  if (argc != 2)
+    {
+      fprintf (stderr, "Usage: %s <hip_module_path>\n", argv[0]);
+      return 1;
+    }
+
+  const auto module_path = argv[1];
+  hipModule_t module;
+  CHECK (hipModuleLoad (&module, module_path));
+
+  /* Launch the kernel.  */
+  hipFunction_t function;
+  CHECK (hipModuleGetFunction (&function, module, "kernel"));
+  CHECK (hipModuleLaunchKernel (function, NUM_BLOCKS, 1, 1,
+				NUM_ITEMS_PER_BLOCK, 1, 1, 0, nullptr, nullptr,
+				nullptr));
+
+  /* Load and unload the module many times.  */
+  for (int i = 0; i < NUM_LOAD_UNLOADS; ++i)
+    {
+      hipModule_t dummy_module;
+      CHECK (hipModuleLoad (&dummy_module, module_path));
+      CHECK (hipModuleUnload (dummy_module));
+    }
+}
+
+#endif
diff --git a/gdb/testsuite/gdb.rocm/code-object-load-while-breakpoint-hit.exp b/gdb/testsuite/gdb.rocm/code-object-load-while-breakpoint-hit.exp
new file mode 100644
index 000000000000..3fe6a952c1d5
--- /dev/null
+++ b/gdb/testsuite/gdb.rocm/code-object-load-while-breakpoint-hit.exp
@@ -0,0 +1,68 @@
+# Copyright 2025 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/>.
+
+# This test verifies what happens when a code object list update happens at the
+# same time as some wave stop events are reported.  It was added following a
+# performance bug fix, where forward progress requirement disabled when
+# pulling events from amd-dbgapi in amd_dbgapi_target_breakpoint::check_status.
+#
+# The test launches a kernel that hits a breakpoint with an always false
+# condition a certain number of times.  Meanwhile, the host loads and unloads
+# a code object in a loop, causing check_status to be called.  The hope is that
+# check_status, when calling process_event_queue, will pull many WAVE_STOP
+# events from the kernel hitting the breakpoint.
+#
+# Without the appropriate fix (of disabling forward progress requirement in
+# check_status), GDB would hit the newly-added assert in process_event_queue,
+# which verifies that forward progress requirement is disabled.  Even without
+# this assert, the test would likely time out (depending on the actual timeout
+# value).
+
+load_lib rocm.exp
+standard_testfile .cpp
+require allow_hipcc_tests
+
+# Build the host executable.
+if { [build_executable "failed to prepare" \
+	  $testfile $srcfile {debug hip}] == -1 } {
+    return -1
+}
+
+set hipmodule_path [standard_output_file ${testfile}.co]
+
+# Build the kernel object file.
+if { [gdb_compile $srcdir/$subdir/$srcfile \
+	$hipmodule_path object \
+	{ debug hip additional_flags=--genco additional_flags=-DDEVICE } ] != "" } {
+    return -1
+}
+
+proc do_test { } {
+    with_rocm_gpu_lock {
+	clean_restart $::binfile
+	gdb_test_no_output "set args $::hipmodule_path" "set args"
+
+	if { ![runto_main] } {
+		return
+	}
+
+	gdb_test "with breakpoint pending on -- break break_here if 0"
+	gdb_continue_to_end "continue to end" "continue" 1
+    }
+}
+
+do_test
-- 
2.49.0

Re: [PATCH v2 3/6] gdb/amd-dbgapi: pass amd_dbgapi_inferior_info to process_event_queue

[Lancelot SIX]

On Mon, Jun 09, 2025 at 12:08:59PM -0400, Simon Marchi wrote:
> New in v2: update a comment
> 
> A following patch will make process_event_queue access a field of
> amd_dbgapi_inferior_info.  Prepare for this by making
> process_event_queue accept an amd_dbgapi_inferior_info object, instead
> of a process id.
> 

This looks good to me, thanks.

Approved-by: Lancelot Six <lancelot.six@amd.com>

Best,
Lancelot.

> Change-Id: I9adc491dd1ff64ff74c40aa7662fffb11bd8332b

Re: [PATCH v2 4/6] gdb/amd-dbgapi: pass amd_dbgapi_inferior_info to process_one_event

[Lancelot SIX]

On Mon, Jun 09, 2025 at 12:09:00PM -0400, Simon Marchi wrote:
> New in v2: this patch is new
> 
> Pass the amd_dbgapi_inferior_info object from process_event_queue to
> process_one_event.  Since process_event_queue pulls events for one
> specific inferior, we know for which inferior the event is.  This
> removes the need for process_one_event to do two dbgapi calls to get the
> relevant pid.  If also removes one inferior lookup.
> 

Hi Simon,

Minor comments below.

> Change-Id: I22927e4b6251513eb3be95785082058aa3d09954
> ---
>  gdb/amd-dbgapi-target.c | 58 +++++++++++++++--------------------------
>  1 file changed, 21 insertions(+), 37 deletions(-)
> 
> diff --git a/gdb/amd-dbgapi-target.c b/gdb/amd-dbgapi-target.c
> index 888552e00d3e..80111359c8cd 100644
> --- a/gdb/amd-dbgapi-target.c
> +++ b/gdb/amd-dbgapi-target.c
> @@ -1143,32 +1143,14 @@ add_gpu_thread (inferior *inf, ptid_t wave_ptid)
>  /* Process an event that was just pulled out of the amd-dbgapi library.  */
>  
>  static void
> -process_one_event (amd_dbgapi_event_id_t event_id,
> +process_one_event (amd_dbgapi_inferior_info &info,
> +		   amd_dbgapi_event_id_t event_id,
>  		   amd_dbgapi_event_kind_t event_kind)
>  {
>    /* Automatically mark this event processed when going out of scope.  */
>    scoped_amd_dbgapi_event_processed mark_event_processed (event_id);
>  
> -  amd_dbgapi_process_id_t process_id;
> -  amd_dbgapi_status_t status
> -    = amd_dbgapi_event_get_info (event_id, AMD_DBGAPI_EVENT_INFO_PROCESS,
> -				 sizeof (process_id), &process_id);
> -  if (status != AMD_DBGAPI_STATUS_SUCCESS)
> -    error (_("event_get_info for event_%ld failed (%s)"), event_id.handle,
> -	   get_status_string (status));
> -
> -  amd_dbgapi_os_process_id_t pid;
> -  status = amd_dbgapi_process_get_info (process_id,
> -					AMD_DBGAPI_PROCESS_INFO_OS_ID,
> -					sizeof (pid), &pid);
> -  if (status != AMD_DBGAPI_STATUS_SUCCESS)
> -    error (_("process_get_info for process_%ld failed (%s)"),
> -	   process_id.handle, get_status_string (status));
> -
> -  auto *proc_target = current_inferior ()->process_target ();
> -  inferior *inf = find_inferior_pid (proc_target, pid);
> -  gdb_assert (inf != nullptr);
> -  amd_dbgapi_inferior_info *info = get_amd_dbgapi_inferior_info (inf);
> +  gdb_assert (info.inf != nullptr);
>  
>    switch (event_kind)
>      {
> @@ -1176,14 +1158,14 @@ process_one_event (amd_dbgapi_event_id_t event_id,
>      case AMD_DBGAPI_EVENT_KIND_WAVE_STOP:
>        {
>  	amd_dbgapi_wave_id_t wave_id;
> -	status
> +	amd_dbgapi_status_t status
>  	  = amd_dbgapi_event_get_info (event_id, AMD_DBGAPI_EVENT_INFO_WAVE,
>  				       sizeof (wave_id), &wave_id);
>  	if (status != AMD_DBGAPI_STATUS_SUCCESS)
>  	  error (_("event_get_info for event_%ld failed (%s)"),
>  		 event_id.handle, get_status_string (status));
>  
> -	ptid_t event_ptid = make_gpu_ptid (pid, wave_id);
> +	ptid_t event_ptid = make_gpu_ptid (info.inf->pid, wave_id);
>  	target_waitstatus ws;
>  
>  	amd_dbgapi_wave_stop_reasons_t stop_reason;
> @@ -1224,9 +1206,10 @@ process_one_event (amd_dbgapi_event_id_t event_id,
>  	    else
>  	      ws.set_stopped (GDB_SIGNAL_0);
>  
> -	    thread_info *thread = proc_target->find_thread (event_ptid);
> -	    if (thread == nullptr)
> -	      thread = add_gpu_thread (inf, event_ptid);
> +            thread_info *thread =
> +                info.inf->process_target()->find_thread(event_ptid);
> +            if (thread == nullptr)

The indentation is off here.  When applying the patch, GIT warns about
indenting with spaces:

    .../binutils-gdb/rebase-apply/patch:69: indent with spaces.
                thread_info *thread =
    .../binutils-gdb/rebase-apply/patch:70: indent with spaces.
                    info.inf->process_target()->find_thread(event_ptid);
    .../binutils-gdb/rebase-apply/patch:71: indent with spaces.
                if (thread == nullptr)
    warning: 3 lines add whitespace errors.

Also, the `=` symbol should be after the newline, not just before it.

With those style issues fixed, this looks good to me, thanks.

Approved-by: Lancelot Six <lancelot.six@amd.com> (amdgpu)

Best,
Lancelot.

> +	      thread = add_gpu_thread (info.inf, event_ptid);
>  
>  	    /* If the wave is stopped because of a software breakpoint, the
>  	       program counter needs to be adjusted so that it points to the
> @@ -1248,7 +1231,7 @@ process_one_event (amd_dbgapi_event_id_t event_id,
>  	  error (_("wave_get_info for wave_%ld failed (%s)"),
>  		 wave_id.handle, get_status_string (status));
>  
> -	info->wave_events.emplace_back (event_ptid, ws);
> +	info.wave_events.emplace_back (event_ptid, ws);
>  	break;
>        }
>  
> @@ -1266,7 +1249,7 @@ process_one_event (amd_dbgapi_event_id_t event_id,
>  	 When amd_dbgapi_target_breakpoint::check_status is called, the current
>  	 inferior is the inferior that hit the breakpoint, which should still be
>  	 the case now.  */
> -      gdb_assert (inf == current_inferior ());
> +      gdb_assert (info.inf == current_inferior ());
>        handle_solib_event ();
>        break;
>  
> @@ -1280,22 +1263,23 @@ process_one_event (amd_dbgapi_event_id_t event_id,
>        {
>  	amd_dbgapi_runtime_state_t runtime_state;
>  
> -	status = amd_dbgapi_event_get_info (event_id,
> -					    AMD_DBGAPI_EVENT_INFO_RUNTIME_STATE,
> -					    sizeof (runtime_state),
> -					    &runtime_state);
> +	amd_dbgapi_status_t status
> +	  = amd_dbgapi_event_get_info (event_id,
> +				       AMD_DBGAPI_EVENT_INFO_RUNTIME_STATE,
> +				       sizeof (runtime_state),
> +				       &runtime_state);
>  	if (status != AMD_DBGAPI_STATUS_SUCCESS)
>  	  error (_("event_get_info for event_%ld failed (%s)"),
>  		 event_id.handle, get_status_string (status));
>  
>  	gdb_assert (runtime_state == AMD_DBGAPI_RUNTIME_STATE_UNLOADED);
>  	gdb_assert
> -	  (info->runtime_state == AMD_DBGAPI_RUNTIME_STATE_LOADED_SUCCESS);
> +	  (info.runtime_state == AMD_DBGAPI_RUNTIME_STATE_LOADED_SUCCESS);
>  
> -	info->runtime_state = runtime_state;
> +	info.runtime_state = runtime_state;
>  
> -	gdb_assert (inf->target_is_pushed (&the_amd_dbgapi_target));
> -	inf->unpush_target (&the_amd_dbgapi_target);
> +	gdb_assert (info.inf->target_is_pushed (&the_amd_dbgapi_target));
> +	info.inf->unpush_target (&the_amd_dbgapi_target);
>        }
>        break;
>  
> @@ -1367,7 +1351,7 @@ process_event_queue (amd_dbgapi_inferior_info &info,
>        if (event_id == AMD_DBGAPI_EVENT_NONE || event_kind == until_event_kind)
>  	return event_id;
>  
> -      process_one_event (event_id, event_kind);
> +      process_one_event (info, event_id, event_kind);
>      }
>  }
>  
> -- 
> 2.49.0

Re: [PATCH v2 4/6] gdb/amd-dbgapi: pass amd_dbgapi_inferior_info to process_one_event

[Simon Marchi]

On 6/16/25 6:32 AM, Lancelot SIX wrote:
> The indentation is off here.  When applying the patch, GIT warns about
> indenting with spaces:
> 
>     .../binutils-gdb/rebase-apply/patch:69: indent with spaces.
>                 thread_info *thread =
>     .../binutils-gdb/rebase-apply/patch:70: indent with spaces.
>                     info.inf->process_target()->find_thread(event_ptid);
>     .../binutils-gdb/rebase-apply/patch:71: indent with spaces.
>                 if (thread == nullptr)
>     warning: 3 lines add whitespace errors.
> 
> Also, the `=` symbol should be after the newline, not just before it.
> 
> With those style issues fixed, this looks good to me, thanks.
> 
> Approved-by: Lancelot Six <lancelot.six@amd.com> (amdgpu)

Thanks, fixed those and pushed the series.

Simon