Add dejagnu tests for cooperative group GWS debugging

[spatrang]

Summary

Add dejagnu coverage for debugging AMD GPU cooperative-group kernels —
i.e. kernels launched via hipLaunchCooperativeKernel /
hipLaunchCooperativeKernelMultiDevice that synchronize at the grid /
multi-grid level. On AMD GPUs these synchronization primitives are
implemented in hardware via Global Wave Sync (GWS), and they have a
distinct wave/scheduling model that has historically only been covered by
out-of-tree tests. This PR brings that coverage into the dejagnu testsuite
so it runs as part of the regular ROCgdb regression suite.

Tests added

File	Scenario
gdb.rocm/coop-group-grid-sync.{cpp,exp}	Single-device cooperative kernel using cooperative_groups::this_grid().sync() (intra-device GWS), launched via hipLaunchCooperativeKernel.
gdb.rocm/coop-group-multi-grid-sync.{cpp,exp}	Multi-device cooperative kernel using both this_grid().sync() and cooperative_groups::this_multi_grid().sync() (intra + cross-device GWS), launched via hipLaunchCooperativeKernelMultiDevice.

What gets verified

coop-group-grid-sync.exp — two sub-tests:

• test_break_around_grid_sync
  • Hit a breakpoint before grid.sync() inside a cooperative dispatch.
  • Confirm multiple AMDGPU Wave threads are stopped (waves participating
    in the GWS barrier).
  • Confirm info dispatches lists the cooperative dispatch.
  • Move the breakpoint to after grid.sync() and continue: it must
    fire (proves GWS-protected code can be debugged across the barrier).
  • Continue to clean program exit.
• test_threads_in_coop_kernel
  • For every AMDGPU Wave parked inside the kernel, switch to it and
    confirm bt 1 reports a frame inside coop_grid_sync_kernel.

coop-group-multi-grid-sync.exp — runs in non-stop mode:

• After continue -a &, confirm a kernel-side breakpoint fires inside
  coop_multi_grid_sync_kernel. Per-GPU child breakpoint instances
  (Breakpoint X.Y) are observed for every participating GPU.
• Continue all threads to program exit, which only succeeds if both
  this_grid().sync() and this_multi_grid().sync() release correctly
  under the debugger.

The host-side post-conditions in the .cpp programs additionally validate
the cooperative semantics numerically (cross-workgroup data dependency for
the single-device case, cross-device sum aggregation for the multi-device
case), so any regression in GWS behavior under the debugger turns into a
test failure rather than a silent miscompare.

Skip / unsupported handling

The tests degrade cleanly on systems that cannot run them:

• Single-device test: queries cooperativeLaunch; if unsupported the
  program prints a recognizable message and exits, and the .exp marks
  the test UNSUPPORTED.
• Multi-device test: requires >= 2 GPUs and
  cooperativeMultiDeviceLaunch on every device. It is also gated by
  the existing hip_devices_support_debug_multi_process requirement.
  Any of those missing → UNSUPPORTED.

No new dejagnu helpers are required; both .exp files use existing
infrastructure in lib/rocm.exp.

Out of scope / follow-ups

Intentionally left out of this PR; happy to extend if reviewers ask:

• Stepping (next / step / stepi) across grid.sync() /
  mgrid.sync() boundaries.
• Conditional breakpoints inside cooperative kernels.
• lane apply / per-lane register inspection while waves are parked at
  the GWS barrier.
• Watchpoints on cooperative shared buffers.

[Copilot]

Pull request overview

Adds new ROCm dejagnu coverage to exercise ROCgdb debugging of cooperative-group HIP kernels that synchronize via GWS, covering both single-device this_grid().sync() and multi-device this_multi_grid().sync() scenarios.

Changes:

• Introduces a single-device cooperative-kernel test that breaks before/after grid.sync() and validates waves/dispatch visibility.
• Introduces a multi-device cooperative-kernel non-stop test that breaks inside a multi-grid kernel and runs through grid + multi-grid barriers to completion.
• Adds two HIP C++ test programs that implement the cooperative-group synchronization patterns and validate results on the host side.

Reviewed changes

Copilot reviewed 4 out of 4 changed files in this pull request and generated 2 comments.

File	Description
gdb/testsuite/gdb.rocm/coop-group-grid-sync.exp	DejaGnu test for single-device cooperative kernel debugging around this_grid().sync().
gdb/testsuite/gdb.rocm/coop-group-grid-sync.cpp	HIP program implementing single-device cooperative grid.sync() and host-side validation.
gdb/testsuite/gdb.rocm/coop-group-multi-grid-sync.exp	DejaGnu non-stop test for multi-device cooperative kernel debugging through this_grid().sync() + this_multi_grid().sync().
gdb/testsuite/gdb.rocm/coop-group-multi-grid-sync.cpp	HIP program implementing multi-device cooperative launch with cross-device aggregation and validation.

---

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

[lancesix]

Hi,
Thanks a lot for this, this is a great starting point.

My main concern for now is gfx110x. We do not support debugging cooperative group on those (documented limitation), the testcase should look for them to not FAIL. This is known that the test will not pass even if the arch do support GWS.

I have added a couple of small comments, I'll get back to a more detailed review after the gfx11 concern has been addressed.

[spatrang]

Hi, Thanks a lot for this, this is a great starting point.

My main concern for now is gfx110x. We do not support debugging cooperative group on those (documented limitation), the testcase should look for them to not FAIL. This is known that the test will not pass even if the arch do support GWS.

I have added a couple of small comments, I'll get back to a more detailed review after the gfx11 concern has been addressed.

Addressed. Added a supports_cooperative_groups helper in lib/rocm.exp that excludes gfx1100/1101/1102/1103, and both .exp files now require it, so on gfx110x the run reports UNSUPPORTED: …: require failed: supports_cooperative_groups instead of FAIL. Mirrors the existing hip_devices_support_debug_multi_process pattern in the same lib.

[aktemur]

We shouldn't need allow-pending anymore. We can optionally use temporary so that we can remove delete_breakpoints below. We can use temporary for the first breakpoint, too.

[spatrang]

Thanks — I applied the temporary part: both breakpoints are now temporary, which let me drop the redundant delete_breakpoints calls (and the same in test_threads_in_coop_kernel).

On allow-pending though - I tried removing it, but it turns out it's still needed here: these breakpoints are on lines inside the kernel (device code), which isn't loaded yet when we set them at main. Without allow-pending, gdb_breakpoint fails outright with "set breakpoint at NN" (gdb declines the unresolved location and defaults to "no" on the pending prompt). On a GPU run this turned into a hard FAIL. So I've kept allow-pending and combined it with temporary (gdb_breakpoint allow-pending temporary). The host-side marker breakpoint in the multi-device test, by contrast, resolves immediately, so temporary alone is fine there.

[aktemur]

Because we stop here before sync, is there a guarantee that we would see 2 distinct workgroups? Wouldn't we have that guarantee rather after synch'ing?

[spatrang]

Good question. The guarantee here comes from the cooperative launch itself rather than from the barrier: hipLaunchCooperativeKernel requires the entire grid to be co-resident on the device for the lifetime of the dispatch (that's precisely what makes grid.sync() safe — a non-co-resident grid could deadlock at the barrier). So all workgroups are resident from dispatch start, including before the first grid.sync(). With just 2 workgroups of 64 threads on the target there's ample occupancy, so both are present. I kept the check pre-sync deliberately: verifying the debugger can see all co-resident waves before the barrier (parked at arbitrary points in the kernel) is a more representative debugging scenario than inspecting them lined up at the sync point. Happy to also add an after-sync check if you'd like the stricter guarantee asserted explicitly.

[aktemur]

Can we use -wrap here, too? It's non-stop mode but only the main thread is supposed to hit the breakpoint. So, I expect we are able to use -wrap and simplify the case to -re "\\\[Inferior 1 \[^\r\n\]* exited normally.*". Please also consider using inferior_exited_re.

[spatrang]

Good suggestion — but this arm went away entirely with the restructure: the marker is now placed on a line reached on every run, so there is no early [Inferior 1 ... exited normally] case to match anymore. No -wrap/inferior_exited_re arm needed here as a result.

[aktemur]

We can do this check early by putting a breakpoint at line 123 and get rid of the "advance to n-gpus-final" check above.

[spatrang]

Done — went with this. Moved the n-gpus-final marker onto the if (n_gpus < N_USED_GPUS) line in the .cpp, which runs on every execution before the inferior's own skip-return. The .exp now does a single gdb_continue_to_breakpoint there, reads n_gpus, and reports unsupported if it's < 2. The dual-arm gdb_test_multiple and its [^\r\n]*\r\n[^\r\n]*\r\n pattern are gone. Validated on gfx942 (in-tree build + 7.14 nightly rocgdb).

[aktemur]

What debugger behavior do we exactly test here? We could put a breakpoint after the sync and all participating blocks/grids would be there. It seems like we are rather testing the runtime, not the debugger.

[spatrang]

Good question — you're right that "do all grids reach the kernel" is a runtime property. The debugger behavior I'm after here is gdb's side: under a single hipLaunchCooperativeKernelMultiDevice dispatch, one source breakpoint resolves to multiple device-side locations, reported as a parent breakpoint with a child instance per GPU (Breakpoint .), and in non-stop mode each device-side stop is observed independently. The loop just confirms gdb reports a stop for every per-device location; the "did every grid arrive" part is left to the host-side result check in the .cpp. I've reworded the in-file comment to make this clearer — happy to switch to the simpler "one breakpoint after the sync" approach if you'd prefer.

[aktemur]

What debugger behavior are we testing here? Before the sync point, stopping waves would be inside the kernel. There is no other kernel. I'm not sure I understand the value of this test from the debugger perspective.

[spatrang]

Thanks - that's a fair point, the bt 1 check was close to tautological with a single kernel. I've adopted approach #1 + #2 to make the debugger value explicit: instead of just confirming each wave's backtrace names the kernel, test_threads_in_coop_kernel now

1. switches to each co-resident wave and reads blockIdx.x, asserting we observe more than one distinct workgroup - i.e. gdb selects the correct per-wave register context; and
2. within one wave, switches between lanes and asserts threadIdx.x differs across lanes - i.e. gdb reports correct per-lane SIMT state. Both are exercised specifically in the co-resident / GWS-barrier context, which is the cooperative-group angle the existing lane/builtin tests don't cover.