gdb/testsuite/gdb.rocm: add multi-inferior stress test

[spatrang]

Summary

Add a stress test that debugs many GPU programs on a single GPU at
once, to exercise the multi-process GPU debug path and GDB's
multi-inferior plumbing under load.

What the test does

• Creates N inferiors of a simple kernel-launcher program with
  add-inferior (not fork), all pinned to the same (default) GPU, and
  runs each to main.
• Plants a breakpoint on the kernel, resumes everything in non-stop
  mode, waits for one kernel stop per inferior, then continues them all
  to a clean exit.
• Running enough programs on one device crosses the per-device limit on
  simultaneously-debugged processes and exercises the firmware context
  switching that kicks in beyond it.

Notes

• Uses add-inferior rather than fork, so the test does not depend on
  fork being available.
• Deliberately runs everything on the default GPU; saturating one
  device is what exercises the per-device limit, and it keeps the test
  shard-friendly (documented in code comments).
• N is chosen to exceed the per-device limit; the test sweeps N/2 and
  N*2 around it. The timeout scales with N, since bringing up many GPU
  runtimes on one device serializes.

Gating

Skipped automatically on targets that do not support multi-process GPU
debugging, via hip_devices_support_debug_multi_process, plus
allow_hip_tests and allow_multi_inferior_tests.

Files added

• gdb/testsuite/gdb.rocm/multi-inferior-stress.cpp
• gdb/testsuite/gdb.rocm/multi-inferior-stress.exp

[lancesix]

I have not re-read that test, but this new test seems to have a lot of similarities with gdb.rocm/multi-inferior-gpu.exp. Couldn't both be merged as a fair amount is common?

[spatrang]

I have not re-read that test, but this new test seems to have a lot of similarities with gdb.rocm/multi-inferior-gpu.exp. Couldn't both be merged as a fair amount is common?

I'd like to keep the two .exp files separate: multi-inferior-gpu.exp is a per-GPU smoke (one child per detected device), while this test deliberately scales N independently of the device count to stress the multi-process GPU debug path and GDB's multi-inferior plumbing. They also have different require gates - the stress test additionally requires allow_multi_inferior_tests. If the duplication concern is the driver loop itself (set non-stop / detach-on-fork / follow-fork → run → wait-for-N hits → continue-each → post-waitpid breakpoint), I'm happy to factor that into a helper in lib/rocm.exp in a follow-up so both .exp files share it.

[lumachad]

I'm in favor of having the common part split out and reused. I think we should keep the tests separate for load-balancing/concurrency purposes.

[lumachad]

@spatrang Could we please have a PR that splits out the common parts that will be reused by this new test, shared with the existing multi-inferior tests? Let's do that first, then we can come back to this one and clear see the differences.

[palves]

I would be better for the long term if this didn't rely on fork. E.g., instead, add inferiors with "add-inferior" and run them.

The reason is that a fork approach won't ever work on Windows. GDB on Windows does not support debugging multiple processes currently, but it's a matter of time until it does. And when it does, we'd want the test to work there too.

Or even, to make it work today on Windows too, we could launch one GDB per inferior. Doable if you juggle spawn_ids. Might give us more coverage if twe exercise both approaches, even. But we could start with one only if it's too complicated.

[lumachad]

I would be better for the long term if this didn't rely on fork. E.g., instead, add inferiors with "add-inferior" and run them.

I think that's a fair requirement, though I think we should put together the testcase as-is for now, and deal with that refactoring at a later stage. Unless we want to adjust the scope of the ticket this is referencing.

[spatrang]

Force-pushed a reworked version addressing the review/JIRA discussion. What changed since the last push:

• No more fork. The test now creates the inferiors with add-inferior + start instead of fork / detach-on-fork off / follow-fork parent, so it doesn't depend on fork being available.
• Single GPU saturation. All inferiors now run on the default GPU instead of being spread across devices with hipSetDevice(idx % num_devices). Saturating one device is what actually exercises the per-device limit and the firmware context switching. This is deliberate and documented in code comments; it also keeps the test shard-friendly.
• N basis. I probed how high N can go on gfx942: no hard cap up to 64 inferiors (no errors), the only cost is wall-clock since bringing up that many GPU runtimes on one device serializes. Given the per-device limit is around 8-16, the test uses N = 16 and sweeps N/2 = 8 and N*2 = 32; 32 is comfortably above the limit. The timeout scales with N.
• The .cpp is now a simple kernel launcher (no parent/child/waitpid).

[czidev-amd]

Once all the other reviewers input are resolved, I have no other objections

[lumachad]

One important point. I think this should go upstream first given it is a generic test, as long as we don't have lack of upstream features to support these tests.

[spatrang]

Thanks for the review @lumachad. Pushed an updated version addressing all the comments (rebased on latest amd-staging):

Cosmetic / comments

• Fixed the block/bracket indentation (the extra with_timeout_factor nesting is gone; it now wraps the do_test call instead).
• Dropped the ;, -- and - connectors from comments; reworded into plain sentences.
• Removed the ticket reference and the redundant gating comment in the header.
• Moved the test description from the .cpp into the .exp; the .cpp header now only describes the program.

Design

• No longer pins to a hardcoded device. The test sets ROCR_VISIBLE_DEVICES to a single GPU, honoring an existing ROCR_VISIBLE_DEVICES (first entry) and otherwise falling back to the first device, applied per-inferior since add-inferior gives each inferior a fresh environment.
• N = 16, and the test now derives and runs N/2 and N*2 automatically.
• Explained the timeout divisor in a comment.
• kern() uses __builtin_amdgcn_s_sleep(1) instead of inline asm.

Verified onROCm 7.14: 166 pass, 0 fail.

On sending this upstream first: agreed it's a generic test with no downstream-only dependency. I'm happy to submit it to upstream GDB first and then let it flow back here. Let me know if you'd like me to proceed that way.

[lumachad]

OK with the comment fixed.